archive-gh-me/wgpu
1
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-12-08 21:26:17 +00:00
Code Issues Projects Releases Wiki Activity

[naga hlsl-out, glsl-out] Support atomicCompareExchangeWeak (#7658)

Browse Source
This commit is contained in:
cryvosh 2025-06-02 07:36:44 -04:00 committed by GitHub
parent 921c6ab597
commit bb46a7f046
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
22 changed files with 1218 additions and 513 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
CHANGELOG.md
View File

@ -55,6 +55,7 @@ Bottom level categories:
- When emitting GLSL, Uniform and Storage Buffer memory layouts are now emitted even if no explicit binding is given. By @cloone8 in [#7579](https://github.com/gfx-rs/wgpu/pull/7579).
- Add support for [quad operations](https://www.w3.org/TR/WGSL/#quad-builtin-functions) (requires `SUBGROUP` feature to be enabled). By @dzamkov and @valaphee in [#7683](https://github.com/gfx-rs/wgpu/pull/7683).
- Add support for `atomicCompareExchangeWeak` in HLSL and GLSL backends. By @cryvosh in [#7658](https://github.com/gfx-rs/wgpu/pull/7658)
### Bug Fixes

87
naga/src/back/glsl/mod.rs
View File

@ -788,6 +788,8 @@ impl<'a, W: Write> Writer<'a, W> {
// you can't make a struct without adding all of its members first.
for (handle, ty) in self.module.types.iter() {
if let TypeInner::Struct { ref members, .. } = ty.inner {
let struct_name = &self.names[&NameKey::Type(handle)];
// Structures ending with runtime-sized arrays can only be
// rendered as shader storage blocks in GLSL, not stand-alone
// struct types.
@ -795,19 +797,19 @@ impl<'a, W: Write> Writer<'a, W> {
.inner
.is_dynamically_sized(&self.module.types)
{
let name = &self.names[&NameKey::Type(handle)];
write!(self.out, "struct {name} ")?;
write!(self.out, "struct {struct_name} ")?;
self.write_struct_body(handle, members)?;
writeln!(self.out, ";")?;
}
}
}
// Write functions to create special types.
// Write functions for special types.
for (type_key, struct_ty) in self.module.special_types.predeclared_types.iter() {
match type_key {
&crate::PredeclaredType::ModfResult { size, scalar }
| &crate::PredeclaredType::FrexpResult { size, scalar } => {
let struct_name = &self.names[&NameKey::Type(*struct_ty)];
let arg_type_name_owner;
let arg_type_name = if let Some(size) = size {
arg_type_name_owner = format!(
@ -836,8 +838,6 @@ impl<'a, W: Write> Writer<'a, W> {
(FREXP_FUNCTION, "frexp", other_type_name)
};
let struct_name = &self.names[&NameKey::Type(*struct_ty)];
writeln!(self.out)?;
if !self.options.version.supports_frexp_function()
&& matches!(type_key, &crate::PredeclaredType::FrexpResult { .. })
@ -861,7 +861,9 @@ impl<'a, W: Write> Writer<'a, W> {
)?;
}
}
&crate::PredeclaredType::AtomicCompareExchangeWeakResult { .. } => {}
&crate::PredeclaredType::AtomicCompareExchangeWeakResult(_) => {
// Handled by the general struct writing loop earlier.
}
}
}
@ -1482,6 +1484,18 @@ impl<'a, W: Write> Writer<'a, W> {
}
}
}
for statement in func.body.iter() {
match *statement {
crate::Statement::Atomic {
fun: crate::AtomicFunction::Exchange { compare: Some(cmp) },
..
} => {
self.need_bake_expressions.insert(cmp);
}
_ => {}
}
}
}
/// Helper method used to get a name for a global
@ -2573,33 +2587,50 @@ impl<'a, W: Write> Writer<'a, W> {
result,
} => {
write!(self.out, "{level}")?;
if let Some(result) = result {
let res_name = Baked(result).to_string();
let res_ty = ctx.resolve_type(result, &self.module.types);
self.write_value_type(res_ty)?;
write!(self.out, " {res_name} = ")?;
self.named_expressions.insert(result, res_name);
}
let fun_str = fun.to_glsl();
write!(self.out, "atomic{fun_str}(")?;
self.write_expr(pointer, ctx)?;
write!(self.out, ", ")?;
// handle the special cases
match *fun {
crate::AtomicFunction::Subtract => {
// we just wrote `InterlockedAdd`, so negate the argument
write!(self.out, "-")?;
crate::AtomicFunction::Exchange {
compare: Some(compare_expr),
} => {
let result_handle = result.expect("CompareExchange must have a result");
let res_name = Baked(result_handle).to_string();
self.write_type(ctx.info[result_handle].ty.handle().unwrap())?;
write!(self.out, " {res_name};")?;
write!(self.out, " {res_name}.old_value = atomicCompSwap(")?;
self.write_expr(pointer, ctx)?;
write!(self.out, ", ")?;
self.write_expr(compare_expr, ctx)?;
write!(self.out, ", ")?;
self.write_expr(value, ctx)?;
writeln!(self.out, ");")?;
write!(
self.out,
"{level}{res_name}.exchanged = ({res_name}.old_value == "
)?;
self.write_expr(compare_expr, ctx)?;
writeln!(self.out, ");")?;
self.named_expressions.insert(result_handle, res_name);
}
crate::AtomicFunction::Exchange { compare: Some(_) } => {
return Err(Error::Custom(
"atomic CompareExchange is not implemented".to_string(),
));
_ => {
if let Some(result) = result {
let res_name = Baked(result).to_string();
self.write_type(ctx.info[result].ty.handle().unwrap())?;
write!(self.out, " {res_name} = ")?;
self.named_expressions.insert(result, res_name);
}
let fun_str = fun.to_glsl();
write!(self.out, "atomic{fun_str}(")?;
self.write_expr(pointer, ctx)?;
write!(self.out, ", ")?;
if let crate::AtomicFunction::Subtract = *fun {
// Emulate `atomicSub` with `atomicAdd` by negating the value.
write!(self.out, "-")?;
}
self.write_expr(value, ctx)?;
writeln!(self.out, ");")?;
}
_ => {}
}
self.write_expr(value, ctx)?;
writeln!(self.out, ");")?;
}
// Stores a value into an image.
Statement::ImageAtomic {

2
naga/src/back/hlsl/conv.rs
View File

@ -222,7 +222,7 @@ impl crate::AtomicFunction {
Self::Min => "Min",
Self::Max => "Max",
Self::Exchange { compare: None } => "Exchange",
Self::Exchange { .. } => "", //TODO
Self::Exchange { .. } => "CompareExchange",
}
}
}

123
naga/src/back/hlsl/writer.rs
View File

@ -269,6 +269,12 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
} => {
self.need_bake_expressions.insert(argument);
}
crate::Statement::Atomic {
fun: crate::AtomicFunction::Exchange { compare: Some(cmp) },
..
} => {
self.need_bake_expressions.insert(cmp);
}
_ => {}
}
}
@ -2358,49 +2364,61 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
result,
} => {
write!(self.out, "{level}")?;
let res_name = match result {
None => None,
Some(result) => {
let name = Baked(result).to_string();
match func_ctx.info[result].ty {
proc::TypeResolution::Handle(handle) => {
self.write_type(module, handle)?
}
proc::TypeResolution::Value(ref value) => {
self.write_value_type(module, value)?
}
};
write!(self.out, " {name}; ")?;
Some((result, name))
}
let res_var_info = if let Some(res_handle) = result {
let name = Baked(res_handle).to_string();
match func_ctx.info[res_handle].ty {
proc::TypeResolution::Handle(handle) => self.write_type(module, handle)?,
proc::TypeResolution::Value(ref value) => {
self.write_value_type(module, value)?
}
};
write!(self.out, " {name}; ")?;
self.named_expressions.insert(res_handle, name.clone());
Some((res_handle, name))
} else {
None
};
// Validation ensures that `pointer` has a `Pointer` type.
let pointer_space = func_ctx
.resolve_type(pointer, &module.types)
.pointer_space()
.unwrap();
let fun_str = fun.to_hlsl_suffix();
let compare_expr = match *fun {
crate::AtomicFunction::Exchange { compare: Some(cmp) } => Some(cmp),
_ => None,
};
match pointer_space {
crate::AddressSpace::WorkGroup => {
write!(self.out, "Interlocked{fun_str}(")?;
self.write_expr(module, pointer, func_ctx)?;
self.emit_hlsl_atomic_tail(
module,
func_ctx,
fun,
compare_expr,
value,
&res_var_info,
)?;
}
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)];
let width = match func_ctx.resolve_type(value, &module.types) {
&TypeInner::Scalar(Scalar { width: 8, .. }) => "64",
_ => "",
};
write!(self.out, "{var_name}.Interlocked{fun_str}{width}(")?;
let chain = mem::take(&mut self.temp_access_chain);
self.write_storage_address(module, &chain, func_ctx)?;
self.temp_access_chain = chain;
self.emit_hlsl_atomic_tail(
module,
func_ctx,
fun,
compare_expr,
value,
&res_var_info,
)?;
}
ref other => {
return Err(Error::Custom(format!(
@ -2408,29 +2426,16 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
)))
}
}
write!(self.out, ", ")?;
// handle the special cases
match *fun {
crate::AtomicFunction::Subtract => {
// we just wrote `InterlockedAdd`, so negate the argument
write!(self.out, "-")?;
if let Some(cmp) = compare_expr {
if let Some(&(_res_handle, ref res_name)) = res_var_info.as_ref() {
write!(
self.out,
"{level}{res_name}.exchanged = ({res_name}.old_value == "
)?;
self.write_expr(module, cmp, func_ctx)?;
writeln!(self.out, ");")?;
}
crate::AtomicFunction::Exchange { compare: Some(_) } => {
return Err(Error::Unimplemented("atomic CompareExchange".to_string()));
}
_ => {}
}
self.write_expr(module, value, func_ctx)?;
// The `original_value` out parameter is optional for all the
// `Interlocked` functions we generate other than
// `InterlockedExchange`.
if let Some((result, name)) = res_name {
write!(self.out, ", {name}")?;
self.named_expressions.insert(result, name);
}
writeln!(self.out, ");")?;
}
Statement::ImageAtomic {
image,
@ -4312,6 +4317,38 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
}
Ok(())
}
/// Helper to emit the shared tail of an HLSL atomic call (arguments, value, result)
fn emit_hlsl_atomic_tail(
&mut self,
module: &Module,
func_ctx: &back::FunctionCtx<'_>,
fun: &crate::AtomicFunction,
compare_expr: Option<Handle<crate::Expression>>,
value: Handle<crate::Expression>,
res_var_info: &Option<(Handle<crate::Expression>, String)>,
) -> BackendResult {
if let Some(cmp) = compare_expr {
write!(self.out, ", ")?;
self.write_expr(module, cmp, func_ctx)?;
}
write!(self.out, ", ")?;
if let crate::AtomicFunction::Subtract = *fun {
// we just wrote `InterlockedAdd`, so negate the argument
write!(self.out, "-")?;
}
self.write_expr(module, value, func_ctx)?;
if let Some(&(_res_handle, ref res_name)) = res_var_info.as_ref() {
write!(self.out, ", ")?;
if compare_expr.is_some() {
write!(self.out, "{res_name}.old_value")?;
} else {
write!(self.out, "{res_name}")?;
}
}
writeln!(self.out, ");")?;
Ok(())
}
}
pub(super) struct MatrixType {

3
naga/tests/in/wgsl/atomicCompareExchange-int64.toml
View File

@ -1,7 +1,8 @@
god_mode = true
targets = "SPIRV | WGSL"
targets = "SPIRV | HLSL | WGSL"
[hlsl]
shader_model = "V6_6"
fake_missing_bindings = true
push_constants_target = { register = 0, space = 0 }
restrict_indexing = true

2
naga/tests/in/wgsl/atomicCompareExchange.toml
View File

@ -1 +1 @@
targets = "SPIRV | METAL | WGSL"
targets = "SPIRV | METAL | GLSL | HLSL | WGSL"

17
naga/tests/in/wgsl/atomicOps-int64.wgsl
View File

@ -129,13 +129,12 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
atomicExchange(&workgroup_struct.atomic_scalar, 1lu);
atomicExchange(&workgroup_struct.atomic_arr[1], 1li);
// // TODO: https://github.com/gpuweb/gpuweb/issues/2021
// atomicCompareExchangeWeak(&storage_atomic_scalar, 1lu);
// atomicCompareExchangeWeak(&storage_atomic_arr[1], 1li);
// atomicCompareExchangeWeak(&storage_struct.atomic_scalar, 1lu);
// atomicCompareExchangeWeak(&storage_struct.atomic_arr[1], 1li);
// atomicCompareExchangeWeak(&workgroup_atomic_scalar, 1lu);
// atomicCompareExchangeWeak(&workgroup_atomic_arr[1], 1li);
// atomicCompareExchangeWeak(&workgroup_struct.atomic_scalar, 1lu);
// atomicCompareExchangeWeak(&workgroup_struct.atomic_arr[1], 1li);
let cas_res_0 = atomicCompareExchangeWeak(&storage_atomic_scalar, 1lu, 2lu);
let cas_res_1 = atomicCompareExchangeWeak(&storage_atomic_arr[1], 1li, 2li);
let cas_res_2 = atomicCompareExchangeWeak(&storage_struct.atomic_scalar, 1lu, 2lu);
let cas_res_3 = atomicCompareExchangeWeak(&storage_struct.atomic_arr[1], 1li, 2li);
let cas_res_4 = atomicCompareExchangeWeak(&workgroup_atomic_scalar, 1lu, 2lu);
let cas_res_5 = atomicCompareExchangeWeak(&workgroup_atomic_arr[1], 1li, 2li);
let cas_res_6 = atomicCompareExchangeWeak(&workgroup_struct.atomic_scalar, 1lu, 2lu);
let cas_res_7 = atomicCompareExchangeWeak(&workgroup_struct.atomic_arr[1], 1li, 2li);
}

17
naga/tests/in/wgsl/atomicOps.wgsl
View File

@ -129,13 +129,12 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
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);
let cas_res_0 = atomicCompareExchangeWeak(&storage_atomic_scalar, 1u, 2u);
let cas_res_1 = atomicCompareExchangeWeak(&storage_atomic_arr[1], 1i, 2i);
let cas_res_2 = atomicCompareExchangeWeak(&storage_struct.atomic_scalar, 1u, 2u);
let cas_res_3 = atomicCompareExchangeWeak(&storage_struct.atomic_arr[1], 1i, 2i);
let cas_res_4 = atomicCompareExchangeWeak(&workgroup_atomic_scalar, 1u, 2u);
let cas_res_5 = atomicCompareExchangeWeak(&workgroup_atomic_arr[1], 1i, 2i);
let cas_res_6 = atomicCompareExchangeWeak(&workgroup_struct.atomic_scalar, 1u, 2u);
let cas_res_7 = atomicCompareExchangeWeak(&workgroup_struct.atomic_arr[1], 1i, 2i);
}

63
naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_i32.Compute.glsl Normal file
View File

@ -0,0 +1,63 @@
#version 310 es
precision highp float;
precision highp int;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
struct _atomic_compare_exchange_resultSint4_ {
int old_value;
bool exchanged;
};
struct _atomic_compare_exchange_resultUint4_ {
uint old_value;
bool exchanged;
};
const uint SIZE = 128u;
layout(std430) buffer type_2_block_0Compute { int _group_0_binding_0_cs[128]; };
void main() {
uint i = 0u;
int old = 0;
bool exchanged = false;
bool loop_init = true;
while(true) {
if (!loop_init) {
uint _e27 = i;
i = (_e27 + 1u);
}
loop_init = false;
uint _e2 = i;
if ((_e2 < SIZE)) {
} else {
break;
}
{
uint _e6 = i;
int _e8 = _group_0_binding_0_cs[_e6];
old = _e8;
exchanged = false;
while(true) {
bool _e12 = exchanged;
if (!(_e12)) {
} else {
break;
}
{
int _e14 = old;
int new = floatBitsToInt((intBitsToFloat(_e14) + 1.0));
uint _e20 = i;
int _e22 = old;
_atomic_compare_exchange_resultSint4_ _e23; _e23.old_value = atomicCompSwap(_group_0_binding_0_cs[_e20], _e22, new);
_e23.exchanged = (_e23.old_value == _e22);
old = _e23.old_value;
exchanged = _e23.exchanged;
}
}
}
}
return;
}

63
naga/tests/out/glsl/wgsl-atomicCompareExchange.test_atomic_compare_exchange_u32.Compute.glsl Normal file
View File

@ -0,0 +1,63 @@
#version 310 es
precision highp float;
precision highp int;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
struct _atomic_compare_exchange_resultSint4_ {
int old_value;
bool exchanged;
};
struct _atomic_compare_exchange_resultUint4_ {
uint old_value;
bool exchanged;
};
const uint SIZE = 128u;
layout(std430) buffer type_4_block_0Compute { uint _group_0_binding_1_cs[128]; };
void main() {
uint i_1 = 0u;
uint old_1 = 0u;
bool exchanged_1 = false;
bool loop_init = true;
while(true) {
if (!loop_init) {
uint _e27 = i_1;
i_1 = (_e27 + 1u);
}
loop_init = false;
uint _e2 = i_1;
if ((_e2 < SIZE)) {
} else {
break;
}
{
uint _e6 = i_1;
uint _e8 = _group_0_binding_1_cs[_e6];
old_1 = _e8;
exchanged_1 = false;
while(true) {
bool _e12 = exchanged_1;
if (!(_e12)) {
} else {
break;
}
{
uint _e14 = old_1;
uint new = floatBitsToUint((uintBitsToFloat(_e14) + 1.0));
uint _e20 = i_1;
uint _e22 = old_1;
_atomic_compare_exchange_resultUint4_ _e23; _e23.old_value = atomicCompSwap(_group_0_binding_1_cs[_e20], _e22, new);
_e23.exchanged = (_e23.old_value == _e22);
old_1 = _e23.old_value;
exchanged_1 = _e23.exchanged;
}
}
}
}
return;
}

24
naga/tests/out/glsl/wgsl-atomicOps.cs_main.Compute.glsl
View File

@ -9,6 +9,14 @@ struct Struct {
uint atomic_scalar;
int atomic_arr[2];
};
struct _atomic_compare_exchange_resultUint4_ {
uint old_value;
bool exchanged;
};
struct _atomic_compare_exchange_resultSint4_ {
int old_value;
bool exchanged;
};
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]; };
@ -127,6 +135,22 @@ void main() {
int _e295 = atomicExchange(workgroup_atomic_arr[1], 1);
uint _e299 = atomicExchange(workgroup_struct.atomic_scalar, 1u);
int _e304 = atomicExchange(workgroup_struct.atomic_arr[1], 1);
_atomic_compare_exchange_resultUint4_ _e308; _e308.old_value = atomicCompSwap(_group_0_binding_0_cs, 1u, 2u);
_e308.exchanged = (_e308.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e313; _e313.old_value = atomicCompSwap(_group_0_binding_1_cs[1], 1, 2);
_e313.exchanged = (_e313.old_value == 1);
_atomic_compare_exchange_resultUint4_ _e318; _e318.old_value = atomicCompSwap(_group_0_binding_2_cs.atomic_scalar, 1u, 2u);
_e318.exchanged = (_e318.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e324; _e324.old_value = atomicCompSwap(_group_0_binding_2_cs.atomic_arr[1], 1, 2);
_e324.exchanged = (_e324.old_value == 1);
_atomic_compare_exchange_resultUint4_ _e328; _e328.old_value = atomicCompSwap(workgroup_atomic_scalar, 1u, 2u);
_e328.exchanged = (_e328.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e333; _e333.old_value = atomicCompSwap(workgroup_atomic_arr[1], 1, 2);
_e333.exchanged = (_e333.old_value == 1);
_atomic_compare_exchange_resultUint4_ _e338; _e338.old_value = atomicCompSwap(workgroup_struct.atomic_scalar, 1u, 2u);
_e338.exchanged = (_e338.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e344; _e344.old_value = atomicCompSwap(workgroup_struct.atomic_arr[1], 1, 2);
_e344.exchanged = (_e344.old_value == 1);
return;
}

127
naga/tests/out/hlsl/wgsl-atomicCompareExchange-int64.hlsl Normal file
View File

@ -0,0 +1,127 @@
struct NagaConstants {
int first_vertex;
int first_instance;
uint other;
};
ConstantBuffer<NagaConstants> _NagaConstants: register(b0, space1);
struct _atomic_compare_exchange_resultSint8_ {
int64_t old_value;
bool exchanged;
int _end_pad_0;
};
struct _atomic_compare_exchange_resultUint8_ {
uint64_t old_value;
bool exchanged;
int _end_pad_0;
};
static const uint SIZE = 128u;
RWByteAddressBuffer arr_i64_ : register(u0);
RWByteAddressBuffer arr_u64_ : register(u1);
[numthreads(1, 1, 1)]
void test_atomic_compare_exchange_i64_()
{
uint i = 0u;
int64_t old = (int64_t)0;
bool exchanged = (bool)0;
uint2 loop_bound = uint2(4294967295u, 4294967295u);
bool loop_init = true;
while(true) {
if (all(loop_bound == uint2(0u, 0u))) { break; }
loop_bound -= uint2(loop_bound.y == 0u, 1u);
if (!loop_init) {
uint _e26 = i;
i = (_e26 + 1u);
}
loop_init = false;
uint _e2 = i;
if ((_e2 < SIZE)) {
} else {
break;
}
{
uint _e6 = i;
int64_t _e8 = arr_i64_.Load<int64_t>(_e6*8);
old = _e8;
exchanged = false;
uint2 loop_bound_1 = uint2(4294967295u, 4294967295u);
while(true) {
if (all(loop_bound_1 == uint2(0u, 0u))) { break; }
loop_bound_1 -= uint2(loop_bound_1.y == 0u, 1u);
bool _e12 = exchanged;
if (!(_e12)) {
} else {
break;
}
{
int64_t _e14 = old;
int64_t new_ = (_e14 + 10L);
uint _e19 = i;
int64_t _e21 = old;
_atomic_compare_exchange_resultSint8_ _e22; arr_i64_.InterlockedCompareExchange64(_e19*8, _e21, new_, _e22.old_value);
_e22.exchanged = (_e22.old_value == _e21);
old = _e22.old_value;
exchanged = _e22.exchanged;
}
}
}
}
return;
}
[numthreads(1, 1, 1)]
void test_atomic_compare_exchange_u64_()
{
uint i_1 = 0u;
uint64_t old_1 = (uint64_t)0;
bool exchanged_1 = (bool)0;
uint2 loop_bound_2 = uint2(4294967295u, 4294967295u);
bool loop_init_1 = true;
while(true) {
if (all(loop_bound_2 == uint2(0u, 0u))) { break; }
loop_bound_2 -= uint2(loop_bound_2.y == 0u, 1u);
if (!loop_init_1) {
uint _e26 = i_1;
i_1 = (_e26 + 1u);
}
loop_init_1 = false;
uint _e2 = i_1;
if ((_e2 < SIZE)) {
} else {
break;
}
{
uint _e6 = i_1;
uint64_t _e8 = arr_u64_.Load<uint64_t>(_e6*8);
old_1 = _e8;
exchanged_1 = false;
uint2 loop_bound_3 = uint2(4294967295u, 4294967295u);
while(true) {
if (all(loop_bound_3 == uint2(0u, 0u))) { break; }
loop_bound_3 -= uint2(loop_bound_3.y == 0u, 1u);
bool _e12 = exchanged_1;
if (!(_e12)) {
} else {
break;
}
{
uint64_t _e14 = old_1;
uint64_t new_1 = (_e14 + 10uL);
uint _e19 = i_1;
uint64_t _e21 = old_1;
_atomic_compare_exchange_resultUint8_ _e22; arr_u64_.InterlockedCompareExchange64(_e19*8, _e21, new_1, _e22.old_value);
_e22.exchanged = (_e22.old_value == _e21);
old_1 = _e22.old_value;
exchanged_1 = _e22.exchanged;
}
}
}
}
return;
}

16
naga/tests/out/hlsl/wgsl-atomicCompareExchange-int64.ron Normal file
View File

@ -0,0 +1,16 @@
(
vertex:[
],
fragment:[
],
compute:[
(
entry_point:"test_atomic_compare_exchange_i64_",
target_profile:"cs_6_6",
),
(
entry_point:"test_atomic_compare_exchange_u64_",
target_profile:"cs_6_6",
),
],
)

118
naga/tests/out/hlsl/wgsl-atomicCompareExchange.hlsl Normal file
View File

@ -0,0 +1,118 @@
struct _atomic_compare_exchange_resultSint4_ {
int old_value;
bool exchanged;
};
struct _atomic_compare_exchange_resultUint4_ {
uint old_value;
bool exchanged;
};
static const uint SIZE = 128u;
RWByteAddressBuffer arr_i32_ : register(u0);
RWByteAddressBuffer arr_u32_ : register(u1);
[numthreads(1, 1, 1)]
void test_atomic_compare_exchange_i32_()
{
uint i = 0u;
int old = (int)0;
bool exchanged = (bool)0;
uint2 loop_bound = uint2(4294967295u, 4294967295u);
bool loop_init = true;
while(true) {
if (all(loop_bound == uint2(0u, 0u))) { break; }
loop_bound -= uint2(loop_bound.y == 0u, 1u);
if (!loop_init) {
uint _e27 = i;
i = (_e27 + 1u);
}
loop_init = false;
uint _e2 = i;
if ((_e2 < SIZE)) {
} else {
break;
}
{
uint _e6 = i;
int _e8 = asint(arr_i32_.Load(_e6*4));
old = _e8;
exchanged = false;
uint2 loop_bound_1 = uint2(4294967295u, 4294967295u);
while(true) {
if (all(loop_bound_1 == uint2(0u, 0u))) { break; }
loop_bound_1 -= uint2(loop_bound_1.y == 0u, 1u);
bool _e12 = exchanged;
if (!(_e12)) {
} else {
break;
}
{
int _e14 = old;
int new_ = asint((asfloat(_e14) + 1.0));
uint _e20 = i;
int _e22 = old;
_atomic_compare_exchange_resultSint4_ _e23; arr_i32_.InterlockedCompareExchange(_e20*4, _e22, new_, _e23.old_value);
_e23.exchanged = (_e23.old_value == _e22);
old = _e23.old_value;
exchanged = _e23.exchanged;
}
}
}
}
return;
}
[numthreads(1, 1, 1)]
void test_atomic_compare_exchange_u32_()
{
uint i_1 = 0u;
uint old_1 = (uint)0;
bool exchanged_1 = (bool)0;
uint2 loop_bound_2 = uint2(4294967295u, 4294967295u);
bool loop_init_1 = true;
while(true) {
if (all(loop_bound_2 == uint2(0u, 0u))) { break; }
loop_bound_2 -= uint2(loop_bound_2.y == 0u, 1u);
if (!loop_init_1) {
uint _e27 = i_1;
i_1 = (_e27 + 1u);
}
loop_init_1 = false;
uint _e2 = i_1;
if ((_e2 < SIZE)) {
} else {
break;
}
{
uint _e6 = i_1;
uint _e8 = asuint(arr_u32_.Load(_e6*4));
old_1 = _e8;
exchanged_1 = false;
uint2 loop_bound_3 = uint2(4294967295u, 4294967295u);
while(true) {
if (all(loop_bound_3 == uint2(0u, 0u))) { break; }
loop_bound_3 -= uint2(loop_bound_3.y == 0u, 1u);
bool _e12 = exchanged_1;
if (!(_e12)) {
} else {
break;
}
{
uint _e14 = old_1;
uint new_1 = asuint((asfloat(_e14) + 1.0));
uint _e20 = i_1;
uint _e22 = old_1;
_atomic_compare_exchange_resultUint4_ _e23; arr_u32_.InterlockedCompareExchange(_e20*4, _e22, new_1, _e23.old_value);
_e23.exchanged = (_e23.old_value == _e22);
old_1 = _e23.old_value;
exchanged_1 = _e23.exchanged;
}
}
}
}
return;
}

16
naga/tests/out/hlsl/wgsl-atomicCompareExchange.ron Normal file
View File

@ -0,0 +1,16 @@
(
vertex:[
],
fragment:[
],
compute:[
(
entry_point:"test_atomic_compare_exchange_i32_",
target_profile:"cs_5_1",
),
(
entry_point:"test_atomic_compare_exchange_u32_",
target_profile:"cs_5_1",
),
],
)

28
naga/tests/out/hlsl/wgsl-atomicOps-int64.hlsl
View File

@ -10,6 +10,18 @@ struct Struct {
int64_t atomic_arr[2];
};
struct _atomic_compare_exchange_resultUint8_ {
uint64_t old_value;
bool exchanged;
int _end_pad_0;
};
struct _atomic_compare_exchange_resultSint8_ {
int64_t old_value;
bool exchanged;
int _end_pad_0;
};
RWByteAddressBuffer storage_atomic_scalar : register(u0);
RWByteAddressBuffer storage_atomic_arr : register(u1);
RWByteAddressBuffer storage_struct : register(u2);
@ -114,5 +126,21 @@ void cs_main(uint3 id : SV_GroupThreadID, uint3 __local_invocation_id : SV_Group
int64_t _e279; InterlockedExchange(workgroup_atomic_arr[1], 1L, _e279);
uint64_t _e283; InterlockedExchange(workgroup_struct.atomic_scalar, 1uL, _e283);
int64_t _e288; InterlockedExchange(workgroup_struct.atomic_arr[1], 1L, _e288);
_atomic_compare_exchange_resultUint8_ _e292; storage_atomic_scalar.InterlockedCompareExchange64(0, 1uL, 2uL, _e292.old_value);
_e292.exchanged = (_e292.old_value == 1uL);
_atomic_compare_exchange_resultSint8_ _e297; storage_atomic_arr.InterlockedCompareExchange64(8, 1L, 2L, _e297.old_value);
_e297.exchanged = (_e297.old_value == 1L);
_atomic_compare_exchange_resultUint8_ _e302; storage_struct.InterlockedCompareExchange64(0, 1uL, 2uL, _e302.old_value);
_e302.exchanged = (_e302.old_value == 1uL);
_atomic_compare_exchange_resultSint8_ _e308; storage_struct.InterlockedCompareExchange64(8+8, 1L, 2L, _e308.old_value);
_e308.exchanged = (_e308.old_value == 1L);
_atomic_compare_exchange_resultUint8_ _e312; InterlockedCompareExchange(workgroup_atomic_scalar, 1uL, 2uL, _e312.old_value);
_e312.exchanged = (_e312.old_value == 1uL);
_atomic_compare_exchange_resultSint8_ _e317; InterlockedCompareExchange(workgroup_atomic_arr[1], 1L, 2L, _e317.old_value);
_e317.exchanged = (_e317.old_value == 1L);
_atomic_compare_exchange_resultUint8_ _e322; InterlockedCompareExchange(workgroup_struct.atomic_scalar, 1uL, 2uL, _e322.old_value);
_e322.exchanged = (_e322.old_value == 1uL);
_atomic_compare_exchange_resultSint8_ _e328; InterlockedCompareExchange(workgroup_struct.atomic_arr[1], 1L, 2L, _e328.old_value);
_e328.exchanged = (_e328.old_value == 1L);
return;
}

26
naga/tests/out/hlsl/wgsl-atomicOps.hlsl
View File

@ -3,6 +3,16 @@ struct Struct {
int atomic_arr[2];
};
struct _atomic_compare_exchange_resultUint4_ {
uint old_value;
bool exchanged;
};
struct _atomic_compare_exchange_resultSint4_ {
int old_value;
bool exchanged;
};
RWByteAddressBuffer storage_atomic_scalar : register(u0);
RWByteAddressBuffer storage_atomic_arr : register(u1);
RWByteAddressBuffer storage_struct : register(u2);
@ -107,5 +117,21 @@ void cs_main(uint3 id : SV_GroupThreadID, uint3 __local_invocation_id : SV_Group
int _e295; InterlockedExchange(workgroup_atomic_arr[1], int(1), _e295);
uint _e299; InterlockedExchange(workgroup_struct.atomic_scalar, 1u, _e299);
int _e304; InterlockedExchange(workgroup_struct.atomic_arr[1], int(1), _e304);
_atomic_compare_exchange_resultUint4_ _e308; storage_atomic_scalar.InterlockedCompareExchange(0, 1u, 2u, _e308.old_value);
_e308.exchanged = (_e308.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e313; storage_atomic_arr.InterlockedCompareExchange(4, int(1), int(2), _e313.old_value);
_e313.exchanged = (_e313.old_value == int(1));
_atomic_compare_exchange_resultUint4_ _e318; storage_struct.InterlockedCompareExchange(0, 1u, 2u, _e318.old_value);
_e318.exchanged = (_e318.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e324; storage_struct.InterlockedCompareExchange(4+4, int(1), int(2), _e324.old_value);
_e324.exchanged = (_e324.old_value == int(1));
_atomic_compare_exchange_resultUint4_ _e328; InterlockedCompareExchange(workgroup_atomic_scalar, 1u, 2u, _e328.old_value);
_e328.exchanged = (_e328.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e333; InterlockedCompareExchange(workgroup_atomic_arr[1], int(1), int(2), _e333.old_value);
_e333.exchanged = (_e333.old_value == int(1));
_atomic_compare_exchange_resultUint4_ _e338; InterlockedCompareExchange(workgroup_struct.atomic_scalar, 1u, 2u, _e338.old_value);
_e338.exchanged = (_e338.old_value == 1u);
_atomic_compare_exchange_resultSint4_ _e344; InterlockedCompareExchange(workgroup_struct.atomic_arr[1], int(1), int(2), _e344.old_value);
_e344.exchanged = (_e344.old_value == int(1));
return;
}

66
naga/tests/out/msl/wgsl-atomicOps.msl
View File

@ -11,6 +11,64 @@ struct Struct {
metal::atomic_uint atomic_scalar;
type_2 atomic_arr;
};
struct _atomic_compare_exchange_resultUint4_ {
uint old_value;
bool exchanged;
};
struct _atomic_compare_exchange_resultSint4_ {
int old_value;
bool exchanged;
};
template <typename A>
_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
device A *atomic_ptr,
uint cmp,
uint v
) {
bool swapped = metal::atomic_compare_exchange_weak_explicit(
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultUint4_{cmp, swapped};
}
template <typename A>
_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
threadgroup A *atomic_ptr,
uint cmp,
uint v
) {
bool swapped = metal::atomic_compare_exchange_weak_explicit(
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultUint4_{cmp, swapped};
}
template <typename A>
_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
device A *atomic_ptr,
int cmp,
int v
) {
bool swapped = metal::atomic_compare_exchange_weak_explicit(
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultSint4_{cmp, swapped};
}
template <typename A>
_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
threadgroup A *atomic_ptr,
int cmp,
int v
) {
bool swapped = metal::atomic_compare_exchange_weak_explicit(
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
);
return _atomic_compare_exchange_resultSint4_{cmp, swapped};
}
struct cs_mainInput {
};
@ -122,5 +180,13 @@ kernel void cs_main(
int _e295 = metal::atomic_exchange_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
uint _e299 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
int _e304 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
_atomic_compare_exchange_resultUint4_ _e308 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_scalar, 1u, 2u);
_atomic_compare_exchange_resultSint4_ _e313 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_arr.inner[1], 1, 2);
_atomic_compare_exchange_resultUint4_ _e318 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_scalar, 1u, 2u);
_atomic_compare_exchange_resultSint4_ _e324 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_arr.inner[1], 1, 2);
_atomic_compare_exchange_resultUint4_ _e328 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_scalar, 1u, 2u);
_atomic_compare_exchange_resultSint4_ _e333 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_arr.inner[1], 1, 2);
_atomic_compare_exchange_resultUint4_ _e338 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_scalar, 1u, 2u);
_atomic_compare_exchange_resultSint4_ _e344 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_arr.inner[1], 1, 2);
return;
}

464
naga/tests/out/spv/wgsl-atomicOps-int64.spvasm
View File

@ -1,31 +1,35 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 193
; Bound: 227
OpCapability Shader
OpCapability Int64Atomics
OpCapability Int64
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %29 "cs_main" %26
OpExecutionMode %29 LocalSize 2 1 1
OpEntryPoint GLCompute %32 "cs_main" %29
OpExecutionMode %32 LocalSize 2 1 1
OpDecorate %5 ArrayStride 8
OpMemberDecorate %8 0 Offset 0
OpMemberDecorate %8 1 Offset 8
OpDecorate %10 DescriptorSet 0
OpDecorate %10 Binding 0
OpDecorate %11 Block
OpMemberDecorate %11 0 Offset 0
OpMemberDecorate %11 1 Offset 8
OpMemberDecorate %12 0 Offset 0
OpMemberDecorate %12 1 Offset 8
OpDecorate %13 DescriptorSet 0
OpDecorate %13 Binding 1
OpDecorate %13 Binding 0
OpDecorate %14 Block
OpMemberDecorate %14 0 Offset 0
OpDecorate %16 DescriptorSet 0
OpDecorate %16 Binding 2
OpDecorate %16 Binding 1
OpDecorate %17 Block
OpMemberDecorate %17 0 Offset 0
OpDecorate %26 BuiltIn LocalInvocationId
OpDecorate %19 DescriptorSet 0
OpDecorate %19 Binding 2
OpDecorate %20 Block
OpMemberDecorate %20 0 Offset 0
OpDecorate %29 BuiltIn LocalInvocationId
%2 = OpTypeVoid
%3 = OpTypeInt 64 0
%4 = OpTypeInt 64 1
@ -34,213 +38,247 @@ OpDecorate %26 BuiltIn LocalInvocationId
%5 = OpTypeArray %4 %6
%8 = OpTypeStruct %3 %5
%9 = OpTypeVector %7 3
%11 = OpTypeStruct %3
%12 = OpTypePointer StorageBuffer %11
%10 = OpVariable %12 StorageBuffer
%14 = OpTypeStruct %5
%10 = OpTypeBool
%11 = OpTypeStruct %3 %10
%12 = OpTypeStruct %4 %10
%14 = OpTypeStruct %3
%15 = OpTypePointer StorageBuffer %14
%13 = OpVariable %15 StorageBuffer
%17 = OpTypeStruct %8
%17 = OpTypeStruct %5
%18 = OpTypePointer StorageBuffer %17
%16 = OpVariable %18 StorageBuffer
%20 = OpTypePointer Workgroup %3
%19 = OpVariable %20 Workgroup
%22 = OpTypePointer Workgroup %5
%21 = OpVariable %22 Workgroup
%24 = OpTypePointer Workgroup %8
%23 = OpVariable %24 Workgroup
%27 = OpTypePointer Input %9
%26 = OpVariable %27 Input
%30 = OpTypeFunction %2
%31 = OpTypePointer StorageBuffer %3
%32 = OpConstant %7 0
%34 = OpTypePointer StorageBuffer %5
%36 = OpTypePointer StorageBuffer %8
%38 = OpConstant %3 1
%39 = OpConstant %4 1
%41 = OpConstantNull %3
%42 = OpConstantNull %5
%43 = OpConstantNull %8
%44 = OpConstantNull %9
%46 = OpTypeBool
%45 = OpTypeVector %46 3
%51 = OpConstant %7 264
%54 = OpTypeInt 32 1
%53 = OpConstant %54 1
%55 = OpConstant %7 64
%56 = OpTypePointer StorageBuffer %4
%57 = OpConstant %7 1
%61 = OpConstant %54 2
%62 = OpConstant %7 256
%63 = OpTypePointer Workgroup %4
%29 = OpFunction %2 None %30
%25 = OpLabel
%28 = OpLoad %9 %26
%33 = OpAccessChain %31 %10 %32
%35 = OpAccessChain %34 %13 %32
%37 = OpAccessChain %36 %16 %32
OpBranch %40
%40 = OpLabel
%47 = OpIEqual %45 %28 %44
%48 = OpAll %46 %47
OpSelectionMerge %49 None
OpBranchConditional %48 %50 %49
%50 = OpLabel
OpStore %19 %41
OpStore %21 %42
OpStore %23 %43
OpBranch %49
%49 = OpLabel
OpControlBarrier %6 %6 %51
OpBranch %52
%52 = OpLabel
OpAtomicStore %33 %53 %55 %38
%58 = OpAccessChain %56 %35 %57
OpAtomicStore %58 %53 %55 %39
%59 = OpAccessChain %31 %37 %32
OpAtomicStore %59 %53 %55 %38
%60 = OpAccessChain %56 %37 %57 %57
OpAtomicStore %60 %53 %55 %39
OpAtomicStore %19 %61 %62 %38
%64 = OpAccessChain %63 %21 %57
OpAtomicStore %64 %61 %62 %39
%65 = OpAccessChain %20 %23 %32
OpAtomicStore %65 %61 %62 %38
%66 = OpAccessChain %63 %23 %57 %57
OpAtomicStore %66 %61 %62 %39
OpControlBarrier %6 %6 %51
%67 = OpAtomicLoad %3 %33 %53 %55
%68 = OpAccessChain %56 %35 %57
%69 = OpAtomicLoad %4 %68 %53 %55
%70 = OpAccessChain %31 %37 %32
%71 = OpAtomicLoad %3 %70 %53 %55
%72 = OpAccessChain %56 %37 %57 %57
%73 = OpAtomicLoad %4 %72 %53 %55
%74 = OpAtomicLoad %3 %19 %61 %62
%75 = OpAccessChain %63 %21 %57
%76 = OpAtomicLoad %4 %75 %61 %62
%77 = OpAccessChain %20 %23 %32
%78 = OpAtomicLoad %3 %77 %61 %62
%79 = OpAccessChain %63 %23 %57 %57
%80 = OpAtomicLoad %4 %79 %61 %62
OpControlBarrier %6 %6 %51
%81 = OpAtomicIAdd %3 %33 %53 %55 %38
%83 = OpAccessChain %56 %35 %57
%82 = OpAtomicIAdd %4 %83 %53 %55 %39
%85 = OpAccessChain %31 %37 %32
%84 = OpAtomicIAdd %3 %85 %53 %55 %38
%87 = OpAccessChain %56 %37 %57 %57
%86 = OpAtomicIAdd %4 %87 %53 %55 %39
%88 = OpAtomicIAdd %3 %19 %61 %62 %38
%90 = OpAccessChain %63 %21 %57
%89 = OpAtomicIAdd %4 %90 %61 %62 %39
%92 = OpAccessChain %20 %23 %32
%91 = OpAtomicIAdd %3 %92 %61 %62 %38
%94 = OpAccessChain %63 %23 %57 %57
%93 = OpAtomicIAdd %4 %94 %61 %62 %39
OpControlBarrier %6 %6 %51
%95 = OpAtomicISub %3 %33 %53 %55 %38
%97 = OpAccessChain %56 %35 %57
%96 = OpAtomicISub %4 %97 %53 %55 %39
%99 = OpAccessChain %31 %37 %32
%98 = OpAtomicISub %3 %99 %53 %55 %38
%101 = OpAccessChain %56 %37 %57 %57
%100 = OpAtomicISub %4 %101 %53 %55 %39
%102 = OpAtomicISub %3 %19 %61 %62 %38
%104 = OpAccessChain %63 %21 %57
%103 = OpAtomicISub %4 %104 %61 %62 %39
%106 = OpAccessChain %20 %23 %32
%105 = OpAtomicISub %3 %106 %61 %62 %38
%108 = OpAccessChain %63 %23 %57 %57
%107 = OpAtomicISub %4 %108 %61 %62 %39
OpControlBarrier %6 %6 %51
%109 = OpAtomicUMax %3 %33 %53 %55 %38
%111 = OpAccessChain %56 %35 %57
%110 = OpAtomicSMax %4 %111 %53 %55 %39
%113 = OpAccessChain %31 %37 %32
%112 = OpAtomicUMax %3 %113 %53 %55 %38
%115 = OpAccessChain %56 %37 %57 %57
%114 = OpAtomicSMax %4 %115 %53 %55 %39
%116 = OpAtomicUMax %3 %19 %61 %62 %38
%118 = OpAccessChain %63 %21 %57
%117 = OpAtomicSMax %4 %118 %61 %62 %39
%120 = OpAccessChain %20 %23 %32
%119 = OpAtomicUMax %3 %120 %61 %62 %38
%122 = OpAccessChain %63 %23 %57 %57
%121 = OpAtomicSMax %4 %122 %61 %62 %39
OpControlBarrier %6 %6 %51
%123 = OpAtomicUMin %3 %33 %53 %55 %38
%125 = OpAccessChain %56 %35 %57
%124 = OpAtomicSMin %4 %125 %53 %55 %39
%127 = OpAccessChain %31 %37 %32
%126 = OpAtomicUMin %3 %127 %53 %55 %38
%129 = OpAccessChain %56 %37 %57 %57
%128 = OpAtomicSMin %4 %129 %53 %55 %39
%130 = OpAtomicUMin %3 %19 %61 %62 %38
%132 = OpAccessChain %63 %21 %57
%131 = OpAtomicSMin %4 %132 %61 %62 %39
%134 = OpAccessChain %20 %23 %32
%133 = OpAtomicUMin %3 %134 %61 %62 %38
%136 = OpAccessChain %63 %23 %57 %57
%135 = OpAtomicSMin %4 %136 %61 %62 %39
OpControlBarrier %6 %6 %51
%137 = OpAtomicAnd %3 %33 %53 %55 %38
%139 = OpAccessChain %56 %35 %57
%138 = OpAtomicAnd %4 %139 %53 %55 %39
%141 = OpAccessChain %31 %37 %32
%140 = OpAtomicAnd %3 %141 %53 %55 %38
%143 = OpAccessChain %56 %37 %57 %57
%142 = OpAtomicAnd %4 %143 %53 %55 %39
%144 = OpAtomicAnd %3 %19 %61 %62 %38
%146 = OpAccessChain %63 %21 %57
%145 = OpAtomicAnd %4 %146 %61 %62 %39
%148 = OpAccessChain %20 %23 %32
%147 = OpAtomicAnd %3 %148 %61 %62 %38
%150 = OpAccessChain %63 %23 %57 %57
%149 = OpAtomicAnd %4 %150 %61 %62 %39
OpControlBarrier %6 %6 %51
%151 = OpAtomicOr %3 %33 %53 %55 %38
%153 = OpAccessChain %56 %35 %57
%152 = OpAtomicOr %4 %153 %53 %55 %39
%155 = OpAccessChain %31 %37 %32
%154 = OpAtomicOr %3 %155 %53 %55 %38
%157 = OpAccessChain %56 %37 %57 %57
%156 = OpAtomicOr %4 %157 %53 %55 %39
%158 = OpAtomicOr %3 %19 %61 %62 %38
%160 = OpAccessChain %63 %21 %57
%159 = OpAtomicOr %4 %160 %61 %62 %39
%162 = OpAccessChain %20 %23 %32
%161 = OpAtomicOr %3 %162 %61 %62 %38
%164 = OpAccessChain %63 %23 %57 %57
%163 = OpAtomicOr %4 %164 %61 %62 %39
OpControlBarrier %6 %6 %51
%165 = OpAtomicXor %3 %33 %53 %55 %38
%167 = OpAccessChain %56 %35 %57
%166 = OpAtomicXor %4 %167 %53 %55 %39
%169 = OpAccessChain %31 %37 %32
%168 = OpAtomicXor %3 %169 %53 %55 %38
%171 = OpAccessChain %56 %37 %57 %57
%170 = OpAtomicXor %4 %171 %53 %55 %39
%172 = OpAtomicXor %3 %19 %61 %62 %38
%174 = OpAccessChain %63 %21 %57
%173 = OpAtomicXor %4 %174 %61 %62 %39
%176 = OpAccessChain %20 %23 %32
%175 = OpAtomicXor %3 %176 %61 %62 %38
%178 = OpAccessChain %63 %23 %57 %57
%177 = OpAtomicXor %4 %178 %61 %62 %39
%179 = OpAtomicExchange %3 %33 %53 %55 %38
%181 = OpAccessChain %56 %35 %57
%180 = OpAtomicExchange %4 %181 %53 %55 %39
%183 = OpAccessChain %31 %37 %32
%182 = OpAtomicExchange %3 %183 %53 %55 %38
%185 = OpAccessChain %56 %37 %57 %57
%184 = OpAtomicExchange %4 %185 %53 %55 %39
%186 = OpAtomicExchange %3 %19 %61 %62 %38
%188 = OpAccessChain %63 %21 %57
%187 = OpAtomicExchange %4 %188 %61 %62 %39
%190 = OpAccessChain %20 %23 %32
%189 = OpAtomicExchange %3 %190 %61 %62 %38
%192 = OpAccessChain %63 %23 %57 %57
%191 = OpAtomicExchange %4 %192 %61 %62 %39
%20 = OpTypeStruct %8
%21 = OpTypePointer StorageBuffer %20
%19 = OpVariable %21 StorageBuffer
%23 = OpTypePointer Workgroup %3
%22 = OpVariable %23 Workgroup
%25 = OpTypePointer Workgroup %5
%24 = OpVariable %25 Workgroup
%27 = OpTypePointer Workgroup %8
%26 = OpVariable %27 Workgroup
%30 = OpTypePointer Input %9
%29 = OpVariable %30 Input
%33 = OpTypeFunction %2
%34 = OpTypePointer StorageBuffer %3
%35 = OpConstant %7 0
%37 = OpTypePointer StorageBuffer %5
%39 = OpTypePointer StorageBuffer %8
%41 = OpConstant %3 1
%42 = OpConstant %4 1
%43 = OpConstant %3 2
%44 = OpConstant %4 2
%46 = OpConstantNull %3
%47 = OpConstantNull %5
%48 = OpConstantNull %8
%49 = OpConstantNull %9
%50 = OpTypeVector %10 3
%55 = OpConstant %7 264
%58 = OpTypeInt 32 1
%57 = OpConstant %58 1
%59 = OpConstant %7 64
%60 = OpTypePointer StorageBuffer %4
%61 = OpConstant %7 1
%65 = OpConstant %58 2
%66 = OpConstant %7 256
%67 = OpTypePointer Workgroup %4
%32 = OpFunction %2 None %33
%28 = OpLabel
%31 = OpLoad %9 %29
%36 = OpAccessChain %34 %13 %35
%38 = OpAccessChain %37 %16 %35
%40 = OpAccessChain %39 %19 %35
OpBranch %45
%45 = OpLabel
%51 = OpIEqual %50 %31 %49
%52 = OpAll %10 %51
OpSelectionMerge %53 None
OpBranchConditional %52 %54 %53
%54 = OpLabel
OpStore %22 %46
OpStore %24 %47
OpStore %26 %48
OpBranch %53
%53 = OpLabel
OpControlBarrier %6 %6 %55
OpBranch %56
%56 = OpLabel
OpAtomicStore %36 %57 %59 %41
%62 = OpAccessChain %60 %38 %61
OpAtomicStore %62 %57 %59 %42
%63 = OpAccessChain %34 %40 %35
OpAtomicStore %63 %57 %59 %41
%64 = OpAccessChain %60 %40 %61 %61
OpAtomicStore %64 %57 %59 %42
OpAtomicStore %22 %65 %66 %41
%68 = OpAccessChain %67 %24 %61
OpAtomicStore %68 %65 %66 %42
%69 = OpAccessChain %23 %26 %35
OpAtomicStore %69 %65 %66 %41
%70 = OpAccessChain %67 %26 %61 %61
OpAtomicStore %70 %65 %66 %42
OpControlBarrier %6 %6 %55
%71 = OpAtomicLoad %3 %36 %57 %59
%72 = OpAccessChain %60 %38 %61
%73 = OpAtomicLoad %4 %72 %57 %59
%74 = OpAccessChain %34 %40 %35
%75 = OpAtomicLoad %3 %74 %57 %59
%76 = OpAccessChain %60 %40 %61 %61
%77 = OpAtomicLoad %4 %76 %57 %59
%78 = OpAtomicLoad %3 %22 %65 %66
%79 = OpAccessChain %67 %24 %61
%80 = OpAtomicLoad %4 %79 %65 %66
%81 = OpAccessChain %23 %26 %35
%82 = OpAtomicLoad %3 %81 %65 %66
%83 = OpAccessChain %67 %26 %61 %61
%84 = OpAtomicLoad %4 %83 %65 %66
OpControlBarrier %6 %6 %55
%85 = OpAtomicIAdd %3 %36 %57 %59 %41
%87 = OpAccessChain %60 %38 %61
%86 = OpAtomicIAdd %4 %87 %57 %59 %42
%89 = OpAccessChain %34 %40 %35
%88 = OpAtomicIAdd %3 %89 %57 %59 %41
%91 = OpAccessChain %60 %40 %61 %61
%90 = OpAtomicIAdd %4 %91 %57 %59 %42
%92 = OpAtomicIAdd %3 %22 %65 %66 %41
%94 = OpAccessChain %67 %24 %61
%93 = OpAtomicIAdd %4 %94 %65 %66 %42
%96 = OpAccessChain %23 %26 %35
%95 = OpAtomicIAdd %3 %96 %65 %66 %41
%98 = OpAccessChain %67 %26 %61 %61
%97 = OpAtomicIAdd %4 %98 %65 %66 %42
OpControlBarrier %6 %6 %55
%99 = OpAtomicISub %3 %36 %57 %59 %41
%101 = OpAccessChain %60 %38 %61
%100 = OpAtomicISub %4 %101 %57 %59 %42
%103 = OpAccessChain %34 %40 %35
%102 = OpAtomicISub %3 %103 %57 %59 %41
%105 = OpAccessChain %60 %40 %61 %61
%104 = OpAtomicISub %4 %105 %57 %59 %42
%106 = OpAtomicISub %3 %22 %65 %66 %41
%108 = OpAccessChain %67 %24 %61
%107 = OpAtomicISub %4 %108 %65 %66 %42
%110 = OpAccessChain %23 %26 %35
%109 = OpAtomicISub %3 %110 %65 %66 %41
%112 = OpAccessChain %67 %26 %61 %61
%111 = OpAtomicISub %4 %112 %65 %66 %42
OpControlBarrier %6 %6 %55
%113 = OpAtomicUMax %3 %36 %57 %59 %41
%115 = OpAccessChain %60 %38 %61
%114 = OpAtomicSMax %4 %115 %57 %59 %42
%117 = OpAccessChain %34 %40 %35
%116 = OpAtomicUMax %3 %117 %57 %59 %41
%119 = OpAccessChain %60 %40 %61 %61
%118 = OpAtomicSMax %4 %119 %57 %59 %42
%120 = OpAtomicUMax %3 %22 %65 %66 %41
%122 = OpAccessChain %67 %24 %61
%121 = OpAtomicSMax %4 %122 %65 %66 %42
%124 = OpAccessChain %23 %26 %35
%123 = OpAtomicUMax %3 %124 %65 %66 %41
%126 = OpAccessChain %67 %26 %61 %61
%125 = OpAtomicSMax %4 %126 %65 %66 %42
OpControlBarrier %6 %6 %55
%127 = OpAtomicUMin %3 %36 %57 %59 %41
%129 = OpAccessChain %60 %38 %61
%128 = OpAtomicSMin %4 %129 %57 %59 %42
%131 = OpAccessChain %34 %40 %35
%130 = OpAtomicUMin %3 %131 %57 %59 %41
%133 = OpAccessChain %60 %40 %61 %61
%132 = OpAtomicSMin %4 %133 %57 %59 %42
%134 = OpAtomicUMin %3 %22 %65 %66 %41
%136 = OpAccessChain %67 %24 %61
%135 = OpAtomicSMin %4 %136 %65 %66 %42
%138 = OpAccessChain %23 %26 %35
%137 = OpAtomicUMin %3 %138 %65 %66 %41
%140 = OpAccessChain %67 %26 %61 %61
%139 = OpAtomicSMin %4 %140 %65 %66 %42
OpControlBarrier %6 %6 %55
%141 = OpAtomicAnd %3 %36 %57 %59 %41
%143 = OpAccessChain %60 %38 %61
%142 = OpAtomicAnd %4 %143 %57 %59 %42
%145 = OpAccessChain %34 %40 %35
%144 = OpAtomicAnd %3 %145 %57 %59 %41
%147 = OpAccessChain %60 %40 %61 %61
%146 = OpAtomicAnd %4 %147 %57 %59 %42
%148 = OpAtomicAnd %3 %22 %65 %66 %41
%150 = OpAccessChain %67 %24 %61
%149 = OpAtomicAnd %4 %150 %65 %66 %42
%152 = OpAccessChain %23 %26 %35
%151 = OpAtomicAnd %3 %152 %65 %66 %41
%154 = OpAccessChain %67 %26 %61 %61
%153 = OpAtomicAnd %4 %154 %65 %66 %42
OpControlBarrier %6 %6 %55
%155 = OpAtomicOr %3 %36 %57 %59 %41
%157 = OpAccessChain %60 %38 %61
%156 = OpAtomicOr %4 %157 %57 %59 %42
%159 = OpAccessChain %34 %40 %35
%158 = OpAtomicOr %3 %159 %57 %59 %41
%161 = OpAccessChain %60 %40 %61 %61
%160 = OpAtomicOr %4 %161 %57 %59 %42
%162 = OpAtomicOr %3 %22 %65 %66 %41
%164 = OpAccessChain %67 %24 %61
%163 = OpAtomicOr %4 %164 %65 %66 %42
%166 = OpAccessChain %23 %26 %35
%165 = OpAtomicOr %3 %166 %65 %66 %41
%168 = OpAccessChain %67 %26 %61 %61
%167 = OpAtomicOr %4 %168 %65 %66 %42
OpControlBarrier %6 %6 %55
%169 = OpAtomicXor %3 %36 %57 %59 %41
%171 = OpAccessChain %60 %38 %61
%170 = OpAtomicXor %4 %171 %57 %59 %42
%173 = OpAccessChain %34 %40 %35
%172 = OpAtomicXor %3 %173 %57 %59 %41
%175 = OpAccessChain %60 %40 %61 %61
%174 = OpAtomicXor %4 %175 %57 %59 %42
%176 = OpAtomicXor %3 %22 %65 %66 %41
%178 = OpAccessChain %67 %24 %61
%177 = OpAtomicXor %4 %178 %65 %66 %42
%180 = OpAccessChain %23 %26 %35
%179 = OpAtomicXor %3 %180 %65 %66 %41
%182 = OpAccessChain %67 %26 %61 %61
%181 = OpAtomicXor %4 %182 %65 %66 %42
%183 = OpAtomicExchange %3 %36 %57 %59 %41
%185 = OpAccessChain %60 %38 %61
%184 = OpAtomicExchange %4 %185 %57 %59 %42
%187 = OpAccessChain %34 %40 %35
%186 = OpAtomicExchange %3 %187 %57 %59 %41
%189 = OpAccessChain %60 %40 %61 %61
%188 = OpAtomicExchange %4 %189 %57 %59 %42
%190 = OpAtomicExchange %3 %22 %65 %66 %41
%192 = OpAccessChain %67 %24 %61
%191 = OpAtomicExchange %4 %192 %65 %66 %42
%194 = OpAccessChain %23 %26 %35
%193 = OpAtomicExchange %3 %194 %65 %66 %41
%196 = OpAccessChain %67 %26 %61 %61
%195 = OpAtomicExchange %4 %196 %65 %66 %42
%198 = OpAtomicCompareExchange %3 %36 %57 %59 %59 %43 %41
%199 = OpIEqual %10 %198 %41
%197 = OpCompositeConstruct %11 %198 %199
%201 = OpAccessChain %60 %38 %61
%202 = OpAtomicCompareExchange %4 %201 %57 %59 %59 %44 %42
%203 = OpIEqual %10 %202 %42
%200 = OpCompositeConstruct %12 %202 %203
%205 = OpAccessChain %34 %40 %35
%206 = OpAtomicCompareExchange %3 %205 %57 %59 %59 %43 %41
%207 = OpIEqual %10 %206 %41
%204 = OpCompositeConstruct %11 %206 %207
%209 = OpAccessChain %60 %40 %61 %61
%210 = OpAtomicCompareExchange %4 %209 %57 %59 %59 %44 %42
%211 = OpIEqual %10 %210 %42
%208 = OpCompositeConstruct %12 %210 %211
%213 = OpAtomicCompareExchange %3 %22 %65 %66 %66 %43 %41
%214 = OpIEqual %10 %213 %41
%212 = OpCompositeConstruct %11 %213 %214
%216 = OpAccessChain %67 %24 %61
%217 = OpAtomicCompareExchange %4 %216 %65 %66 %66 %44 %42
%218 = OpIEqual %10 %217 %42
%215 = OpCompositeConstruct %12 %217 %218
%220 = OpAccessChain %23 %26 %35
%221 = OpAtomicCompareExchange %3 %220 %65 %66 %66 %43 %41
%222 = OpIEqual %10 %221 %41
%219 = OpCompositeConstruct %11 %221 %222
%224 = OpAccessChain %67 %26 %61 %61
%225 = OpAtomicCompareExchange %4 %224 %65 %66 %66 %44 %42
%226 = OpIEqual %10 %225 %42
%223 = OpCompositeConstruct %12 %225 %226
OpReturn
OpFunctionEnd

452
naga/tests/out/spv/wgsl-atomicOps.spvasm
View File

@ -1,29 +1,33 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 189
; Bound: 221
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %28 "cs_main" %25
OpExecutionMode %28 LocalSize 2 1 1
OpEntryPoint GLCompute %31 "cs_main" %28
OpExecutionMode %31 LocalSize 2 1 1
OpDecorate %5 ArrayStride 4
OpMemberDecorate %7 0 Offset 0
OpMemberDecorate %7 1 Offset 4
OpDecorate %9 DescriptorSet 0
OpDecorate %9 Binding 0
OpDecorate %10 Block
OpMemberDecorate %10 0 Offset 0
OpMemberDecorate %10 1 Offset 4
OpMemberDecorate %11 0 Offset 0
OpMemberDecorate %11 1 Offset 4
OpDecorate %12 DescriptorSet 0
OpDecorate %12 Binding 1
OpDecorate %12 Binding 0
OpDecorate %13 Block
OpMemberDecorate %13 0 Offset 0
OpDecorate %15 DescriptorSet 0
OpDecorate %15 Binding 2
OpDecorate %15 Binding 1
OpDecorate %16 Block
OpMemberDecorate %16 0 Offset 0
OpDecorate %25 BuiltIn LocalInvocationId
OpDecorate %18 DescriptorSet 0
OpDecorate %18 Binding 2
OpDecorate %19 Block
OpMemberDecorate %19 0 Offset 0
OpDecorate %28 BuiltIn LocalInvocationId
%2 = OpTypeVoid
%3 = OpTypeInt 32 0
%4 = OpTypeInt 32 1
@ -31,210 +35,242 @@ OpDecorate %25 BuiltIn LocalInvocationId
%5 = OpTypeArray %4 %6
%7 = OpTypeStruct %3 %5
%8 = OpTypeVector %3 3
%10 = OpTypeStruct %3
%11 = OpTypePointer StorageBuffer %10
%9 = OpVariable %11 StorageBuffer
%13 = OpTypeStruct %5
%9 = OpTypeBool
%10 = OpTypeStruct %3 %9
%11 = OpTypeStruct %4 %9
%13 = OpTypeStruct %3
%14 = OpTypePointer StorageBuffer %13
%12 = OpVariable %14 StorageBuffer
%16 = OpTypeStruct %7
%16 = OpTypeStruct %5
%17 = OpTypePointer StorageBuffer %16
%15 = OpVariable %17 StorageBuffer
%19 = OpTypePointer Workgroup %3
%18 = OpVariable %19 Workgroup
%21 = OpTypePointer Workgroup %5
%20 = OpVariable %21 Workgroup
%23 = OpTypePointer Workgroup %7
%22 = OpVariable %23 Workgroup
%26 = OpTypePointer Input %8
%25 = OpVariable %26 Input
%29 = OpTypeFunction %2
%30 = OpTypePointer StorageBuffer %3
%31 = OpConstant %3 0
%33 = OpTypePointer StorageBuffer %5
%35 = OpTypePointer StorageBuffer %7
%37 = OpConstant %3 1
%38 = OpConstant %4 1
%40 = OpConstantNull %3
%41 = OpConstantNull %5
%42 = OpConstantNull %7
%43 = OpConstantNull %8
%45 = OpTypeBool
%44 = OpTypeVector %45 3
%50 = OpConstant %3 264
%52 = OpConstant %3 64
%53 = OpTypePointer StorageBuffer %4
%57 = OpConstant %4 2
%58 = OpConstant %3 256
%59 = OpTypePointer Workgroup %4
%28 = OpFunction %2 None %29
%24 = OpLabel
%27 = OpLoad %8 %25
%32 = OpAccessChain %30 %9 %31
%34 = OpAccessChain %33 %12 %31
%36 = OpAccessChain %35 %15 %31
OpBranch %39
%39 = OpLabel
%46 = OpIEqual %44 %27 %43
%47 = OpAll %45 %46
OpSelectionMerge %48 None
OpBranchConditional %47 %49 %48
%49 = OpLabel
OpStore %18 %40
OpStore %20 %41
OpStore %22 %42
OpBranch %48
%48 = OpLabel
OpControlBarrier %6 %6 %50
%19 = OpTypeStruct %7
%20 = OpTypePointer StorageBuffer %19
%18 = OpVariable %20 StorageBuffer
%22 = OpTypePointer Workgroup %3
%21 = OpVariable %22 Workgroup
%24 = OpTypePointer Workgroup %5
%23 = OpVariable %24 Workgroup
%26 = OpTypePointer Workgroup %7
%25 = OpVariable %26 Workgroup
%29 = OpTypePointer Input %8
%28 = OpVariable %29 Input
%32 = OpTypeFunction %2
%33 = OpTypePointer StorageBuffer %3
%34 = OpConstant %3 0
%36 = OpTypePointer StorageBuffer %5
%38 = OpTypePointer StorageBuffer %7
%40 = OpConstant %3 1
%41 = OpConstant %4 1
%42 = OpConstant %4 2
%44 = OpConstantNull %3
%45 = OpConstantNull %5
%46 = OpConstantNull %7
%47 = OpConstantNull %8
%48 = OpTypeVector %9 3
%53 = OpConstant %3 264
%55 = OpConstant %3 64
%56 = OpTypePointer StorageBuffer %4
%60 = OpConstant %3 256
%61 = OpTypePointer Workgroup %4
%31 = OpFunction %2 None %32
%27 = OpLabel
%30 = OpLoad %8 %28
%35 = OpAccessChain %33 %12 %34
%37 = OpAccessChain %36 %15 %34
%39 = OpAccessChain %38 %18 %34
OpBranch %43
%43 = OpLabel
%49 = OpIEqual %48 %30 %47
%50 = OpAll %9 %49
OpSelectionMerge %51 None
OpBranchConditional %50 %52 %51
%52 = OpLabel
OpStore %21 %44
OpStore %23 %45
OpStore %25 %46
OpBranch %51
%51 = OpLabel
OpAtomicStore %32 %38 %52 %37
%54 = OpAccessChain %53 %34 %37
OpAtomicStore %54 %38 %52 %38
%55 = OpAccessChain %30 %36 %31
OpAtomicStore %55 %38 %52 %37
%56 = OpAccessChain %53 %36 %37 %37
OpAtomicStore %56 %38 %52 %38
OpAtomicStore %18 %57 %58 %37
%60 = OpAccessChain %59 %20 %37
OpAtomicStore %60 %57 %58 %38
%61 = OpAccessChain %19 %22 %31
OpAtomicStore %61 %57 %58 %37
%62 = OpAccessChain %59 %22 %37 %37
OpAtomicStore %62 %57 %58 %38
OpControlBarrier %6 %6 %50
%63 = OpAtomicLoad %3 %32 %38 %52
%64 = OpAccessChain %53 %34 %37
%65 = OpAtomicLoad %4 %64 %38 %52
%66 = OpAccessChain %30 %36 %31
%67 = OpAtomicLoad %3 %66 %38 %52
%68 = OpAccessChain %53 %36 %37 %37
%69 = OpAtomicLoad %4 %68 %38 %52
%70 = OpAtomicLoad %3 %18 %57 %58
%71 = OpAccessChain %59 %20 %37
%72 = OpAtomicLoad %4 %71 %57 %58
%73 = OpAccessChain %19 %22 %31
%74 = OpAtomicLoad %3 %73 %57 %58
%75 = OpAccessChain %59 %22 %37 %37
%76 = OpAtomicLoad %4 %75 %57 %58
OpControlBarrier %6 %6 %50
%77 = OpAtomicIAdd %3 %32 %38 %52 %37
%79 = OpAccessChain %53 %34 %37
%78 = OpAtomicIAdd %4 %79 %38 %52 %38
%81 = OpAccessChain %30 %36 %31
%80 = OpAtomicIAdd %3 %81 %38 %52 %37
%83 = OpAccessChain %53 %36 %37 %37
%82 = OpAtomicIAdd %4 %83 %38 %52 %38
%84 = OpAtomicIAdd %3 %18 %57 %58 %37
%86 = OpAccessChain %59 %20 %37
%85 = OpAtomicIAdd %4 %86 %57 %58 %38
%88 = OpAccessChain %19 %22 %31
%87 = OpAtomicIAdd %3 %88 %57 %58 %37
%90 = OpAccessChain %59 %22 %37 %37
%89 = OpAtomicIAdd %4 %90 %57 %58 %38
OpControlBarrier %6 %6 %50
%91 = OpAtomicISub %3 %32 %38 %52 %37
%93 = OpAccessChain %53 %34 %37
%92 = OpAtomicISub %4 %93 %38 %52 %38
%95 = OpAccessChain %30 %36 %31
%94 = OpAtomicISub %3 %95 %38 %52 %37
%97 = OpAccessChain %53 %36 %37 %37
%96 = OpAtomicISub %4 %97 %38 %52 %38
%98 = OpAtomicISub %3 %18 %57 %58 %37
%100 = OpAccessChain %59 %20 %37
%99 = OpAtomicISub %4 %100 %57 %58 %38
%102 = OpAccessChain %19 %22 %31
%101 = OpAtomicISub %3 %102 %57 %58 %37
%104 = OpAccessChain %59 %22 %37 %37
%103 = OpAtomicISub %4 %104 %57 %58 %38
OpControlBarrier %6 %6 %50
%105 = OpAtomicUMax %3 %32 %38 %52 %37
%107 = OpAccessChain %53 %34 %37
%106 = OpAtomicSMax %4 %107 %38 %52 %38
%109 = OpAccessChain %30 %36 %31
%108 = OpAtomicUMax %3 %109 %38 %52 %37
%111 = OpAccessChain %53 %36 %37 %37
%110 = OpAtomicSMax %4 %111 %38 %52 %38
%112 = OpAtomicUMax %3 %18 %57 %58 %37
%114 = OpAccessChain %59 %20 %37
%113 = OpAtomicSMax %4 %114 %57 %58 %38
%116 = OpAccessChain %19 %22 %31
%115 = OpAtomicUMax %3 %116 %57 %58 %37
%118 = OpAccessChain %59 %22 %37 %37
%117 = OpAtomicSMax %4 %118 %57 %58 %38
OpControlBarrier %6 %6 %50
%119 = OpAtomicUMin %3 %32 %38 %52 %37
%121 = OpAccessChain %53 %34 %37
%120 = OpAtomicSMin %4 %121 %38 %52 %38
%123 = OpAccessChain %30 %36 %31
%122 = OpAtomicUMin %3 %123 %38 %52 %37
%125 = OpAccessChain %53 %36 %37 %37
%124 = OpAtomicSMin %4 %125 %38 %52 %38
%126 = OpAtomicUMin %3 %18 %57 %58 %37
%128 = OpAccessChain %59 %20 %37
%127 = OpAtomicSMin %4 %128 %57 %58 %38
%130 = OpAccessChain %19 %22 %31
%129 = OpAtomicUMin %3 %130 %57 %58 %37
%132 = OpAccessChain %59 %22 %37 %37
%131 = OpAtomicSMin %4 %132 %57 %58 %38
OpControlBarrier %6 %6 %50
%133 = OpAtomicAnd %3 %32 %38 %52 %37
%135 = OpAccessChain %53 %34 %37
%134 = OpAtomicAnd %4 %135 %38 %52 %38
%137 = OpAccessChain %30 %36 %31
%136 = OpAtomicAnd %3 %137 %38 %52 %37
%139 = OpAccessChain %53 %36 %37 %37
%138 = OpAtomicAnd %4 %139 %38 %52 %38
%140 = OpAtomicAnd %3 %18 %57 %58 %37
%142 = OpAccessChain %59 %20 %37
%141 = OpAtomicAnd %4 %142 %57 %58 %38
%144 = OpAccessChain %19 %22 %31
%143 = OpAtomicAnd %3 %144 %57 %58 %37
%146 = OpAccessChain %59 %22 %37 %37
%145 = OpAtomicAnd %4 %146 %57 %58 %38
OpControlBarrier %6 %6 %50
%147 = OpAtomicOr %3 %32 %38 %52 %37
%149 = OpAccessChain %53 %34 %37
%148 = OpAtomicOr %4 %149 %38 %52 %38
%151 = OpAccessChain %30 %36 %31
%150 = OpAtomicOr %3 %151 %38 %52 %37
%153 = OpAccessChain %53 %36 %37 %37
%152 = OpAtomicOr %4 %153 %38 %52 %38
%154 = OpAtomicOr %3 %18 %57 %58 %37
%156 = OpAccessChain %59 %20 %37
%155 = OpAtomicOr %4 %156 %57 %58 %38
%158 = OpAccessChain %19 %22 %31
%157 = OpAtomicOr %3 %158 %57 %58 %37
%160 = OpAccessChain %59 %22 %37 %37
%159 = OpAtomicOr %4 %160 %57 %58 %38
OpControlBarrier %6 %6 %50
%161 = OpAtomicXor %3 %32 %38 %52 %37
%163 = OpAccessChain %53 %34 %37
%162 = OpAtomicXor %4 %163 %38 %52 %38
%165 = OpAccessChain %30 %36 %31
%164 = OpAtomicXor %3 %165 %38 %52 %37
%167 = OpAccessChain %53 %36 %37 %37
%166 = OpAtomicXor %4 %167 %38 %52 %38
%168 = OpAtomicXor %3 %18 %57 %58 %37
%170 = OpAccessChain %59 %20 %37
%169 = OpAtomicXor %4 %170 %57 %58 %38
%172 = OpAccessChain %19 %22 %31
%171 = OpAtomicXor %3 %172 %57 %58 %37
%174 = OpAccessChain %59 %22 %37 %37
%173 = OpAtomicXor %4 %174 %57 %58 %38
%175 = OpAtomicExchange %3 %32 %38 %52 %37
%177 = OpAccessChain %53 %34 %37
%176 = OpAtomicExchange %4 %177 %38 %52 %38
%179 = OpAccessChain %30 %36 %31
%178 = OpAtomicExchange %3 %179 %38 %52 %37
%181 = OpAccessChain %53 %36 %37 %37
%180 = OpAtomicExchange %4 %181 %38 %52 %38
%182 = OpAtomicExchange %3 %18 %57 %58 %37
%184 = OpAccessChain %59 %20 %37
%183 = OpAtomicExchange %4 %184 %57 %58 %38
%186 = OpAccessChain %19 %22 %31
%185 = OpAtomicExchange %3 %186 %57 %58 %37
%188 = OpAccessChain %59 %22 %37 %37
%187 = OpAtomicExchange %4 %188 %57 %58 %38
OpControlBarrier %6 %6 %53
OpBranch %54
%54 = OpLabel
OpAtomicStore %35 %41 %55 %40
%57 = OpAccessChain %56 %37 %40
OpAtomicStore %57 %41 %55 %41
%58 = OpAccessChain %33 %39 %34
OpAtomicStore %58 %41 %55 %40
%59 = OpAccessChain %56 %39 %40 %40
OpAtomicStore %59 %41 %55 %41
OpAtomicStore %21 %42 %60 %40
%62 = OpAccessChain %61 %23 %40
OpAtomicStore %62 %42 %60 %41
%63 = OpAccessChain %22 %25 %34
OpAtomicStore %63 %42 %60 %40
%64 = OpAccessChain %61 %25 %40 %40
OpAtomicStore %64 %42 %60 %41
OpControlBarrier %6 %6 %53
%65 = OpAtomicLoad %3 %35 %41 %55
%66 = OpAccessChain %56 %37 %40
%67 = OpAtomicLoad %4 %66 %41 %55
%68 = OpAccessChain %33 %39 %34
%69 = OpAtomicLoad %3 %68 %41 %55
%70 = OpAccessChain %56 %39 %40 %40
%71 = OpAtomicLoad %4 %70 %41 %55
%72 = OpAtomicLoad %3 %21 %42 %60
%73 = OpAccessChain %61 %23 %40
%74 = OpAtomicLoad %4 %73 %42 %60
%75 = OpAccessChain %22 %25 %34
%76 = OpAtomicLoad %3 %75 %42 %60
%77 = OpAccessChain %61 %25 %40 %40
%78 = OpAtomicLoad %4 %77 %42 %60
OpControlBarrier %6 %6 %53
%79 = OpAtomicIAdd %3 %35 %41 %55 %40
%81 = OpAccessChain %56 %37 %40
%80 = OpAtomicIAdd %4 %81 %41 %55 %41
%83 = OpAccessChain %33 %39 %34
%82 = OpAtomicIAdd %3 %83 %41 %55 %40
%85 = OpAccessChain %56 %39 %40 %40
%84 = OpAtomicIAdd %4 %85 %41 %55 %41
%86 = OpAtomicIAdd %3 %21 %42 %60 %40
%88 = OpAccessChain %61 %23 %40
%87 = OpAtomicIAdd %4 %88 %42 %60 %41
%90 = OpAccessChain %22 %25 %34
%89 = OpAtomicIAdd %3 %90 %42 %60 %40
%92 = OpAccessChain %61 %25 %40 %40
%91 = OpAtomicIAdd %4 %92 %42 %60 %41
OpControlBarrier %6 %6 %53
%93 = OpAtomicISub %3 %35 %41 %55 %40
%95 = OpAccessChain %56 %37 %40
%94 = OpAtomicISub %4 %95 %41 %55 %41
%97 = OpAccessChain %33 %39 %34
%96 = OpAtomicISub %3 %97 %41 %55 %40
%99 = OpAccessChain %56 %39 %40 %40
%98 = OpAtomicISub %4 %99 %41 %55 %41
%100 = OpAtomicISub %3 %21 %42 %60 %40
%102 = OpAccessChain %61 %23 %40
%101 = OpAtomicISub %4 %102 %42 %60 %41
%104 = OpAccessChain %22 %25 %34
%103 = OpAtomicISub %3 %104 %42 %60 %40
%106 = OpAccessChain %61 %25 %40 %40
%105 = OpAtomicISub %4 %106 %42 %60 %41
OpControlBarrier %6 %6 %53
%107 = OpAtomicUMax %3 %35 %41 %55 %40
%109 = OpAccessChain %56 %37 %40
%108 = OpAtomicSMax %4 %109 %41 %55 %41
%111 = OpAccessChain %33 %39 %34
%110 = OpAtomicUMax %3 %111 %41 %55 %40
%113 = OpAccessChain %56 %39 %40 %40
%112 = OpAtomicSMax %4 %113 %41 %55 %41
%114 = OpAtomicUMax %3 %21 %42 %60 %40
%116 = OpAccessChain %61 %23 %40
%115 = OpAtomicSMax %4 %116 %42 %60 %41
%118 = OpAccessChain %22 %25 %34
%117 = OpAtomicUMax %3 %118 %42 %60 %40
%120 = OpAccessChain %61 %25 %40 %40
%119 = OpAtomicSMax %4 %120 %42 %60 %41
OpControlBarrier %6 %6 %53
%121 = OpAtomicUMin %3 %35 %41 %55 %40
%123 = OpAccessChain %56 %37 %40
%122 = OpAtomicSMin %4 %123 %41 %55 %41
%125 = OpAccessChain %33 %39 %34
%124 = OpAtomicUMin %3 %125 %41 %55 %40
%127 = OpAccessChain %56 %39 %40 %40
%126 = OpAtomicSMin %4 %127 %41 %55 %41
%128 = OpAtomicUMin %3 %21 %42 %60 %40
%130 = OpAccessChain %61 %23 %40
%129 = OpAtomicSMin %4 %130 %42 %60 %41
%132 = OpAccessChain %22 %25 %34
%131 = OpAtomicUMin %3 %132 %42 %60 %40
%134 = OpAccessChain %61 %25 %40 %40
%133 = OpAtomicSMin %4 %134 %42 %60 %41
OpControlBarrier %6 %6 %53
%135 = OpAtomicAnd %3 %35 %41 %55 %40
%137 = OpAccessChain %56 %37 %40
%136 = OpAtomicAnd %4 %137 %41 %55 %41
%139 = OpAccessChain %33 %39 %34
%138 = OpAtomicAnd %3 %139 %41 %55 %40
%141 = OpAccessChain %56 %39 %40 %40
%140 = OpAtomicAnd %4 %141 %41 %55 %41
%142 = OpAtomicAnd %3 %21 %42 %60 %40
%144 = OpAccessChain %61 %23 %40
%143 = OpAtomicAnd %4 %144 %42 %60 %41
%146 = OpAccessChain %22 %25 %34
%145 = OpAtomicAnd %3 %146 %42 %60 %40
%148 = OpAccessChain %61 %25 %40 %40
%147 = OpAtomicAnd %4 %148 %42 %60 %41
OpControlBarrier %6 %6 %53
%149 = OpAtomicOr %3 %35 %41 %55 %40
%151 = OpAccessChain %56 %37 %40
%150 = OpAtomicOr %4 %151 %41 %55 %41
%153 = OpAccessChain %33 %39 %34
%152 = OpAtomicOr %3 %153 %41 %55 %40
%155 = OpAccessChain %56 %39 %40 %40
%154 = OpAtomicOr %4 %155 %41 %55 %41
%156 = OpAtomicOr %3 %21 %42 %60 %40
%158 = OpAccessChain %61 %23 %40
%157 = OpAtomicOr %4 %158 %42 %60 %41
%160 = OpAccessChain %22 %25 %34
%159 = OpAtomicOr %3 %160 %42 %60 %40
%162 = OpAccessChain %61 %25 %40 %40
%161 = OpAtomicOr %4 %162 %42 %60 %41
OpControlBarrier %6 %6 %53
%163 = OpAtomicXor %3 %35 %41 %55 %40
%165 = OpAccessChain %56 %37 %40
%164 = OpAtomicXor %4 %165 %41 %55 %41
%167 = OpAccessChain %33 %39 %34
%166 = OpAtomicXor %3 %167 %41 %55 %40
%169 = OpAccessChain %56 %39 %40 %40
%168 = OpAtomicXor %4 %169 %41 %55 %41
%170 = OpAtomicXor %3 %21 %42 %60 %40
%172 = OpAccessChain %61 %23 %40
%171 = OpAtomicXor %4 %172 %42 %60 %41
%174 = OpAccessChain %22 %25 %34
%173 = OpAtomicXor %3 %174 %42 %60 %40
%176 = OpAccessChain %61 %25 %40 %40
%175 = OpAtomicXor %4 %176 %42 %60 %41
%177 = OpAtomicExchange %3 %35 %41 %55 %40
%179 = OpAccessChain %56 %37 %40
%178 = OpAtomicExchange %4 %179 %41 %55 %41
%181 = OpAccessChain %33 %39 %34
%180 = OpAtomicExchange %3 %181 %41 %55 %40
%183 = OpAccessChain %56 %39 %40 %40
%182 = OpAtomicExchange %4 %183 %41 %55 %41
%184 = OpAtomicExchange %3 %21 %42 %60 %40
%186 = OpAccessChain %61 %23 %40
%185 = OpAtomicExchange %4 %186 %42 %60 %41
%188 = OpAccessChain %22 %25 %34
%187 = OpAtomicExchange %3 %188 %42 %60 %40
%190 = OpAccessChain %61 %25 %40 %40
%189 = OpAtomicExchange %4 %190 %42 %60 %41
%192 = OpAtomicCompareExchange %3 %35 %41 %55 %55 %6 %40
%193 = OpIEqual %9 %192 %40
%191 = OpCompositeConstruct %10 %192 %193
%195 = OpAccessChain %56 %37 %40
%196 = OpAtomicCompareExchange %4 %195 %41 %55 %55 %42 %41
%197 = OpIEqual %9 %196 %41
%194 = OpCompositeConstruct %11 %196 %197
%199 = OpAccessChain %33 %39 %34
%200 = OpAtomicCompareExchange %3 %199 %41 %55 %55 %6 %40
%201 = OpIEqual %9 %200 %40
%198 = OpCompositeConstruct %10 %200 %201
%203 = OpAccessChain %56 %39 %40 %40
%204 = OpAtomicCompareExchange %4 %203 %41 %55 %55 %42 %41
%205 = OpIEqual %9 %204 %41
%202 = OpCompositeConstruct %11 %204 %205
%207 = OpAtomicCompareExchange %3 %21 %42 %60 %60 %6 %40
%208 = OpIEqual %9 %207 %40
%206 = OpCompositeConstruct %10 %207 %208
%210 = OpAccessChain %61 %23 %40
%211 = OpAtomicCompareExchange %4 %210 %42 %60 %60 %42 %41
%212 = OpIEqual %9 %211 %41
%209 = OpCompositeConstruct %11 %211 %212
%214 = OpAccessChain %22 %25 %34
%215 = OpAtomicCompareExchange %3 %214 %42 %60 %60 %6 %40
%216 = OpIEqual %9 %215 %40
%213 = OpCompositeConstruct %10 %215 %216
%218 = OpAccessChain %61 %25 %40 %40
%219 = OpAtomicCompareExchange %4 %218 %42 %60 %60 %42 %41
%220 = OpIEqual %9 %219 %41
%217 = OpCompositeConstruct %11 %219 %220
OpReturn
OpFunctionEnd

8
naga/tests/out/wgsl/wgsl-atomicOps-int64.wgsl
View File

@ -103,5 +103,13 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
let _e279 = atomicExchange((&workgroup_atomic_arr[1]), 1li);
let _e283 = atomicExchange((&workgroup_struct.atomic_scalar), 1lu);
let _e288 = atomicExchange((&workgroup_struct.atomic_arr[1]), 1li);
let _e292 = atomicCompareExchangeWeak((&storage_atomic_scalar), 1lu, 2lu);
let _e297 = atomicCompareExchangeWeak((&storage_atomic_arr[1]), 1li, 2li);
let _e302 = atomicCompareExchangeWeak((&storage_struct.atomic_scalar), 1lu, 2lu);
let _e308 = atomicCompareExchangeWeak((&storage_struct.atomic_arr[1]), 1li, 2li);
let _e312 = atomicCompareExchangeWeak((&workgroup_atomic_scalar), 1lu, 2lu);
let _e317 = atomicCompareExchangeWeak((&workgroup_atomic_arr[1]), 1li, 2li);
let _e322 = atomicCompareExchangeWeak((&workgroup_struct.atomic_scalar), 1lu, 2lu);
let _e328 = atomicCompareExchangeWeak((&workgroup_struct.atomic_arr[1]), 1li, 2li);
return;
}

8
naga/tests/out/wgsl/wgsl-atomicOps.wgsl
View File

@ -103,5 +103,13 @@ fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
let _e295 = atomicExchange((&workgroup_atomic_arr[1]), 1i);
let _e299 = atomicExchange((&workgroup_struct.atomic_scalar), 1u);
let _e304 = atomicExchange((&workgroup_struct.atomic_arr[1]), 1i);
let _e308 = atomicCompareExchangeWeak((&storage_atomic_scalar), 1u, 2u);
let _e313 = atomicCompareExchangeWeak((&storage_atomic_arr[1]), 1i, 2i);
let _e318 = atomicCompareExchangeWeak((&storage_struct.atomic_scalar), 1u, 2u);
let _e324 = atomicCompareExchangeWeak((&storage_struct.atomic_arr[1]), 1i, 2i);
let _e328 = atomicCompareExchangeWeak((&workgroup_atomic_scalar), 1u, 2u);
let _e333 = atomicCompareExchangeWeak((&workgroup_atomic_arr[1]), 1i, 2i);
let _e338 = atomicCompareExchangeWeak((&workgroup_struct.atomic_scalar), 1u, 2u);
let _e344 = atomicCompareExchangeWeak((&workgroup_struct.atomic_arr[1]), 1i, 2i);
return;
}