nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-02-27 14:23:40 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add atomicSubtract support

Browse Source
This commit is contained in:
Dzmitry Malyshau 2021-08-24 17:45:29 -04:00 committed by Dzmitry Malyshau
parent 715a53f891
commit 1e9f2b9287
14 changed files with 96 additions and 45 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

20
src/back/glsl/mod.rs
View File

@ -71,7 +71,7 @@ pub type BindingMap = std::collections::BTreeMap<crate::ResourceBinding, u8>;
impl crate::AtomicFunction {
fn to_glsl(self) -> &'static str {
match self {
Self::Add => "Add",
Self::Add | Self::Subtract => "Add",
Self::And => "And",
Self::InclusiveOr => "Or",
Self::ExclusiveOr => "Xor",
@ -1706,12 +1706,20 @@ impl<'a, W: Write> Writer<'a, W> {
let fun_str = fun.to_glsl();
write!(self.out, "atomic{}(", fun_str)?;
self.write_expr(pointer, ctx)?;
if let crate::AtomicFunction::Exchange { compare: Some(_) } = *fun {
return Err(Error::Custom(
"atomic CompareExchange is not implemented".to_string(),
));
}
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(_) } => {
return Err(Error::Custom(
"atomic CompareExchange is not implemented".to_string(),
));
}
_ => {}
}
self.write_expr(value, ctx)?;
writeln!(self.out, ");")?;
}

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

@ -133,7 +133,7 @@ impl crate::AtomicFunction {
/// Return the HLSL suffix for the `InterlockedXxx` method.
pub(super) fn to_hlsl_suffix(self) -> &'static str {
match self {
Self::Add => "Add",
Self::Add | Self::Subtract => "Add",
Self::And => "And",
Self::InclusiveOr => "Or",
Self::ExclusiveOr => "Xor",

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

@ -1374,10 +1374,18 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
res_name, var_name, fun_str
)?;
self.write_storage_address(module, &chain, func_ctx)?;
if let crate::AtomicFunction::Exchange { compare: Some(_) } = *fun {
return Err(Error::Unimplemented("atomic CompareExchange".to_string()));
}
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(_) } => {
return Err(Error::Unimplemented("atomic CompareExchange".to_string()));
}
_ => {}
}
self.write_expr(module, value, func_ctx)?;
writeln!(self.out, ", {});", res_name)?;
self.temp_access_chain = chain;

7
src/back/msl/writer.rs
View File

@ -1647,6 +1647,9 @@ impl<W: Write> Writer<W> {
crate::AtomicFunction::Add => {
self.put_atomic_fetch(pointer, "add", value, &context.expression)?;
}
crate::AtomicFunction::Subtract => {
self.put_atomic_fetch(pointer, "sub", value, &context.expression)?;
}
crate::AtomicFunction::And => {
self.put_atomic_fetch(pointer, "and", value, &context.expression)?;
}
@ -2643,8 +2646,8 @@ fn test_stack_size() {
}
let stack_size = addresses.end - addresses.start;
// check the size (in debug only)
// last observed macOS value: 17504
if !(13000..=19000).contains(&stack_size) {
// last observed macOS value: 19152 (CI)
if !(13000..=20000).contains(&stack_size) {
panic!("`put_block` stack size {} has changed!", stack_size);
}
}

9
src/back/spv/block.rs
View File

@ -1262,6 +1262,15 @@ impl<'w> BlockContext<'w> {
semantics_id,
value_id,
),
crate::AtomicFunction::Subtract => Instruction::atomic_binary(
spirv::Op::AtomicISub,
result_type_id,
id,
pointer_id,
scope_constant_id,
semantics_id,
value_id,
),
crate::AtomicFunction::And => Instruction::atomic_binary(
spirv::Op::AtomicAnd,
result_type_id,

1
src/back/wgsl/mod.rs
View File

@ -31,6 +31,7 @@ impl crate::AtomicFunction {
fn to_wgsl(self) -> &'static str {
match self {
Self::Add => "Add",
Self::Subtract => "Sub",
Self::And => "And",
Self::InclusiveOr => "Or",
Self::ExclusiveOr => "Xor",

9
src/front/wgsl/mod.rs
View File

@ -1403,6 +1403,15 @@ impl Parser {
)?;
return Ok(Some(handle));
}
"atomicSub" => {
let _ = lexer.next();
let handle = self.parse_atomic_helper(
lexer,
crate::AtomicFunction::Subtract,
ctx.reborrow(),
)?;
return Ok(Some(handle));
}
"atomicAnd" => {
let _ = lexer.next();
let handle = self.parse_atomic_helper(

1
src/lib.rs
View File

@ -750,6 +750,7 @@ pub enum BinaryOperator {
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
pub enum AtomicFunction {
Add,
Subtract,
And,
ExclusiveOr,
InclusiveOr,

1
tests/in/access.wgsl
View File

@ -43,6 +43,7 @@ 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);

14
tests/out/glsl/access.atomics.Compute.glsl
View File

@ -18,18 +18,20 @@ void main() {
int value = _group_0_binding_0.atom;
int _e6 = atomicAdd(_group_0_binding_0.atom, 5);
tmp = _e6;
int _e9 = atomicAnd(_group_0_binding_0.atom, 5);
int _e9 = atomicAdd(_group_0_binding_0.atom, -5);
tmp = _e9;
int _e12 = atomicOr(_group_0_binding_0.atom, 5);
int _e12 = atomicAnd(_group_0_binding_0.atom, 5);
tmp = _e12;
int _e15 = atomicXor(_group_0_binding_0.atom, 5);
int _e15 = atomicOr(_group_0_binding_0.atom, 5);
tmp = _e15;
int _e18 = atomicMin(_group_0_binding_0.atom, 5);
int _e18 = atomicXor(_group_0_binding_0.atom, 5);
tmp = _e18;
int _e21 = atomicMax(_group_0_binding_0.atom, 5);
int _e21 = atomicMin(_group_0_binding_0.atom, 5);
tmp = _e21;
int _e24 = atomicExchange(_group_0_binding_0.atom, 5);
int _e24 = atomicMax(_group_0_binding_0.atom, 5);
tmp = _e24;
int _e27 = atomicExchange(_group_0_binding_0.atom, 5);
tmp = _e27;
_group_0_binding_0.atom = value;
return;
}

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

@ -50,18 +50,20 @@ void atomics()
int value = asint(bar.Load(64));
int _e6; bar.InterlockedAdd(64, 5, _e6);
tmp = _e6;
int _e9; bar.InterlockedAnd(64, 5, _e9);
int _e9; bar.InterlockedAdd(64, -5, _e9);
tmp = _e9;
int _e12; bar.InterlockedOr(64, 5, _e12);
int _e12; bar.InterlockedAnd(64, 5, _e12);
tmp = _e12;
int _e15; bar.InterlockedXor(64, 5, _e15);
int _e15; bar.InterlockedOr(64, 5, _e15);
tmp = _e15;
int _e18; bar.InterlockedMin(64, 5, _e18);
int _e18; bar.InterlockedXor(64, 5, _e18);
tmp = _e18;
int _e21; bar.InterlockedMax(64, 5, _e21);
int _e21; bar.InterlockedMin(64, 5, _e21);
tmp = _e21;
int _e24; bar.InterlockedExchange(64, 5, _e24);
int _e24; bar.InterlockedMax(64, 5, _e24);
tmp = _e24;
int _e27; bar.InterlockedExchange(64, 5, _e27);
tmp = _e27;
bar.Store(64, asuint(value));
return;
}

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

@ -58,18 +58,20 @@ kernel void atomics(
int value = metal::atomic_load_explicit(&bar.atom, metal::memory_order_relaxed);
int _e6 = metal::atomic_fetch_add_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e6;
int _e9 = metal::atomic_fetch_and_explicit(&bar.atom, 5, metal::memory_order_relaxed);
int _e9 = metal::atomic_fetch_sub_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e9;
int _e12 = metal::atomic_fetch_or_explicit(&bar.atom, 5, metal::memory_order_relaxed);
int _e12 = metal::atomic_fetch_and_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e12;
int _e15 = metal::atomic_fetch_xor_explicit(&bar.atom, 5, metal::memory_order_relaxed);
int _e15 = metal::atomic_fetch_or_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e15;
int _e18 = metal::atomic_fetch_min_explicit(&bar.atom, 5, metal::memory_order_relaxed);
int _e18 = metal::atomic_fetch_xor_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e18;
int _e21 = metal::atomic_fetch_max_explicit(&bar.atom, 5, metal::memory_order_relaxed);
int _e21 = metal::atomic_fetch_min_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e21;
int _e24 = metal::atomic_exchange_explicit(&bar.atom, 5, metal::memory_order_relaxed);
int _e24 = metal::atomic_fetch_max_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e24;
int _e27 = metal::atomic_exchange_explicit(&bar.atom, 5, metal::memory_order_relaxed);
tmp = _e27;
metal::atomic_store_explicit(&bar.atom, value, metal::memory_order_relaxed);
return;
}

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

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 105
; Bound: 107
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
@ -139,24 +139,27 @@ OpBranch %85
%90 = OpAtomicIAdd %4 %91 %11 %89 %16
OpStore %82 %90
%93 = OpAccessChain %86 %27 %15
%92 = OpAtomicAnd %4 %93 %11 %89 %16
%92 = OpAtomicISub %4 %93 %11 %89 %16
OpStore %82 %92
%95 = OpAccessChain %86 %27 %15
%94 = OpAtomicOr %4 %95 %11 %89 %16
%94 = OpAtomicAnd %4 %95 %11 %89 %16
OpStore %82 %94
%97 = OpAccessChain %86 %27 %15
%96 = OpAtomicXor %4 %97 %11 %89 %16
%96 = OpAtomicOr %4 %97 %11 %89 %16
OpStore %82 %96
%99 = OpAccessChain %86 %27 %15
%98 = OpAtomicSMin %4 %99 %11 %89 %16
%98 = OpAtomicXor %4 %99 %11 %89 %16
OpStore %82 %98
%101 = OpAccessChain %86 %27 %15
%100 = OpAtomicSMax %4 %101 %11 %89 %16
%100 = OpAtomicSMin %4 %101 %11 %89 %16
OpStore %82 %100
%103 = OpAccessChain %86 %27 %15
%102 = OpAtomicExchange %4 %103 %11 %89 %16
%102 = OpAtomicSMax %4 %103 %11 %89 %16
OpStore %82 %102
%104 = OpAccessChain %86 %27 %15
OpAtomicStore %104 %11 %89 %88
%105 = OpAccessChain %86 %27 %15
%104 = OpAtomicExchange %4 %105 %11 %89 %16
OpStore %82 %104
%106 = OpAccessChain %86 %27 %15
OpAtomicStore %106 %11 %89 %88
OpReturn
OpFunctionEnd

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

@ -37,18 +37,20 @@ fn atomics() {
let value: i32 = atomicLoad(&bar.atom);
let _e6: i32 = atomicAdd(&bar.atom, 5);
tmp = _e6;
let _e9: i32 = atomicAnd(&bar.atom, 5);
let _e9: i32 = atomicSub(&bar.atom, 5);
tmp = _e9;
let _e12: i32 = atomicOr(&bar.atom, 5);
let _e12: i32 = atomicAnd(&bar.atom, 5);
tmp = _e12;
let _e15: i32 = atomicXor(&bar.atom, 5);
let _e15: i32 = atomicOr(&bar.atom, 5);
tmp = _e15;
let _e18: i32 = atomicMin(&bar.atom, 5);
let _e18: i32 = atomicXor(&bar.atom, 5);
tmp = _e18;
let _e21: i32 = atomicMax(&bar.atom, 5);
let _e21: i32 = atomicMin(&bar.atom, 5);
tmp = _e21;
let _e24: i32 = atomicExchange(&bar.atom, 5);
let _e24: i32 = atomicMax(&bar.atom, 5);
tmp = _e24;
let _e27: i32 = atomicExchange(&bar.atom, 5);
tmp = _e27;
atomicStore(&bar.atom, value);
return;
}