nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-11-25 00:03:29 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add support for 64 bit integer atomic operations in shaders.

Browse Source 
Add the following flags to `wgpu_types::Features`:

- `SHADER_INT64_ATOMIC_ALL_OPS` enables all atomic operations on `atomic<i64>` and
  `atomic<u64>` values.

- `SHADER_INT64_ATOMIC_MIN_MAX` is a subset of the above, enabling only
  `AtomicFunction::Min` and `AtomicFunction::Max` operations on `atomic<i64>` and
  `atomic<u64>` values in the `Storage` address space. These are the only 64-bit
  atomic operations available on Metal as of 3.1.

Add corresponding flags to `naga::valid::Capabilities`. These are supported by the
WGSL front end, and all Naga backends.

Platform support:

- On Direct3d 12, in `D3D12_FEATURE_DATA_D3D12_OPTIONS9`, if
  `AtomicInt64OnTypedResourceSupported` and `AtomicInt64OnGroupSharedSupported` are
  both available, then both wgpu features described above are available.

- On Metal, `SHADER_INT64_ATOMIC_MIN_MAX` is available on Apple9 hardware, and on
  hardware that advertises both Apple8 and Mac2 support. This also requires Metal
  Shading Language 2.4 or later. Metal does not yet support the more general
  `SHADER_INT64_ATOMIC_ALL_OPS`.

- On Vulkan, if the `VK_KHR_shader_atomic_int64` extension is available with both the
  `shader_buffer_int64_atomics` and `shader_shared_int64_atomics` features, then both
  wgpu features described above are available.
This commit is contained in:
Atlas Dostal 2024-06-08 17:21:25 -07:00 committed by Jim Blandy
parent 583cc6ab04
commit abba12ae4e
51 changed files with 1916 additions and 160 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

35
CHANGELOG.md
View File

@ -81,8 +81,41 @@ for message in compilation_info
By @stefnotch in [#5410](https://github.com/gfx-rs/wgpu/pull/5410) By @stefnotch in [#5410](https://github.com/gfx-rs/wgpu/pull/5410)
### New features #### 64 bit integer atomic support in shaders.
Add support for 64 bit integer atomic operations in shaders.
Add the following flags to `wgpu_types::Features`:
- `SHADER_INT64_ATOMIC_ALL_OPS` enables all atomic operations on `atomic<i64>` and
`atomic<u64>` values.
- `SHADER_INT64_ATOMIC_MIN_MAX` is a subset of the above, enabling only
`AtomicFunction::Min` and `AtomicFunction::Max` operations on `atomic<i64>` and
`atomic<u64>` values in the `Storage` address space. These are the only 64-bit
atomic operations available on Metal as of 3.1.
Add corresponding flags to `naga::valid::Capabilities`. These are supported by the
WGSL front end, and all Naga backends.
Platform support:
- On Direct3d 12, in `D3D12_FEATURE_DATA_D3D12_OPTIONS9`, if
`AtomicInt64OnTypedResourceSupported` and `AtomicInt64OnGroupSharedSupported` are
both available, then both wgpu features described above are available.
- On Metal, `SHADER_INT64_ATOMIC_MIN_MAX` is available on Apple9 hardware, and on
hardware that advertises both Apple8 and Mac2 support. This also requires Metal
Shading Language 2.4 or later. Metal does not yet support the more general
`SHADER_INT64_ATOMIC_ALL_OPS`.
- On Vulkan, if the `VK_KHR_shader_atomic_int64` extension is available with both the
`shader_buffer_int64_atomics` and `shader_shared_int64_atomics` features, then both
wgpu features described above are available.
By @atlv24 in [#5383](https://github.com/gfx-rs/wgpu/pull/5383)
### New features
#### Vulkan #### Vulkan
- Added a `PipelineCache` resource to allow using Vulkan pipeline caches. By @DJMcNab in [#5319](https://github.com/gfx-rs/wgpu/pull/5319) - Added a `PipelineCache` resource to allow using Vulkan pipeline caches. By @DJMcNab in [#5319](https://github.com/gfx-rs/wgpu/pull/5319)

4
naga/src/back/dot/mod.rs
View File

@ -244,7 +244,9 @@ impl StatementGraph {
value, value,
result, result,
} => { } => {
self.emits.push((id, result)); if let Some(result) = result {
self.emits.push((id, result));
}
self.dependencies.push((id, pointer, "pointer")); self.dependencies.push((id, pointer, "pointer"));
self.dependencies.push((id, value, "value")); self.dependencies.push((id, value, "value"));
if let crate::AtomicFunction::Exchange { compare: Some(cmp) } = *fun { if let crate::AtomicFunction::Exchange { compare: Some(cmp) } = *fun {

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

@ -2368,11 +2368,13 @@ impl<'a, W: Write> Writer<'a, W> {
result, result,
} => { } => {
write!(self.out, "{level}")?; write!(self.out, "{level}")?;
let res_name = format!("{}{}", back::BAKE_PREFIX, result.index()); if let Some(result) = result {
let res_ty = ctx.resolve_type(result, &self.module.types); let res_name = format!("{}{}", back::BAKE_PREFIX, result.index());
self.write_value_type(res_ty)?; let res_ty = ctx.resolve_type(result, &self.module.types);
write!(self.out, " {res_name} = ")?; self.write_value_type(res_ty)?;
self.named_expressions.insert(result, res_name); write!(self.out, " {res_name} = ")?;
self.named_expressions.insert(result, res_name);
}
let fun_str = fun.to_glsl(); let fun_str = fun.to_glsl();
write!(self.out, "atomic{fun_str}(")?; write!(self.out, "atomic{fun_str}(")?;

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

@ -1919,11 +1919,20 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
result, result,
} => { } => {
write!(self.out, "{level}")?; write!(self.out, "{level}")?;
let res_name = format!("{}{}", back::BAKE_PREFIX, result.index()); let res_name = match result {
match func_ctx.info[result].ty { None => None,
proc::TypeResolution::Handle(handle) => self.write_type(module, handle)?, Some(result) => {
proc::TypeResolution::Value(ref value) => { let name = format!("{}{}", back::BAKE_PREFIX, result.index());
self.write_value_type(module, value)? 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))
} }
}; };
@ -1934,7 +1943,6 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
.unwrap(); .unwrap();
let fun_str = fun.to_hlsl_suffix(); let fun_str = fun.to_hlsl_suffix();
write!(self.out, " {res_name}; ")?;
match pointer_space { match pointer_space {
crate::AddressSpace::WorkGroup => { crate::AddressSpace::WorkGroup => {
write!(self.out, "Interlocked{fun_str}(")?; write!(self.out, "Interlocked{fun_str}(")?;
@ -1970,8 +1978,16 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
_ => {} _ => {}
} }
self.write_expr(module, value, func_ctx)?; self.write_expr(module, value, func_ctx)?;
writeln!(self.out, ", {res_name});")?;
self.named_expressions.insert(result, res_name); // 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::WorkGroupUniformLoad { pointer, result } => { Statement::WorkGroupUniformLoad { pointer, result } => {
self.write_barrier(crate::Barrier::WORK_GROUP, level)?; self.write_barrier(crate::Barrier::WORK_GROUP, level)?;

23
naga/src/back/msl/mod.rs
View File

@ -1,6 +1,9 @@
/*! /*!
Backend for [MSL][msl] (Metal Shading Language). Backend for [MSL][msl] (Metal Shading Language).
This backend does not support the [`SHADER_INT64_ATOMIC_ALL_OPS`][all-atom]
capability.
## Binding model ## Binding model
Metal's bindings are flat per resource. Since there isn't an obvious mapping Metal's bindings are flat per resource. Since there isn't an obvious mapping
@ -24,6 +27,8 @@ For the result type, if it's a structure, we re-compose it with a temporary valu
holding the result. holding the result.
[msl]: https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf [msl]: https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
[all-atom]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_ALL_OPS
*/ */
use crate::{arena::Handle, proc::index, valid::ModuleInfo}; use crate::{arena::Handle, proc::index, valid::ModuleInfo};
@ -661,21 +666,3 @@ fn test_error_size() {
use std::mem::size_of; use std::mem::size_of;
assert_eq!(size_of::<Error>(), 32); assert_eq!(size_of::<Error>(), 32);
} }
impl crate::AtomicFunction {
fn to_msl(self) -> Result<&'static str, Error> {
Ok(match self {
Self::Add => "fetch_add",
Self::Subtract => "fetch_sub",
Self::And => "fetch_and",
Self::InclusiveOr => "fetch_or",
Self::ExclusiveOr => "fetch_xor",
Self::Min => "fetch_min",
Self::Max => "fetch_max",
Self::Exchange { compare: None } => "exchange",
Self::Exchange { compare: Some(_) } => Err(Error::FeatureNotImplemented(
"atomic CompareExchange".to_string(),
))?,
})
}
}

47
naga/src/back/msl/writer.rs
View File

@ -3058,11 +3058,22 @@ impl<W: Write> Writer<W> {
value, value,
result, result,
} => { } => {
// This backend supports `SHADER_INT64_ATOMIC_MIN_MAX` but not
// `SHADER_INT64_ATOMIC_ALL_OPS`, so we can assume that if `result` is
// `Some`, we are not operating on a 64-bit value, and that if we are
// operating on a 64-bit value, `result` is `None`.
write!(self.out, "{level}")?; write!(self.out, "{level}")?;
let res_name = format!("{}{}", back::BAKE_PREFIX, result.index()); let fun_str = if let Some(result) = result {
self.start_baking_expression(result, &context.expression, &res_name)?; let res_name = format!("{}{}", back::BAKE_PREFIX, result.index());
self.named_expressions.insert(result, res_name); self.start_baking_expression(result, &context.expression, &res_name)?;
let fun_str = fun.to_msl()?; self.named_expressions.insert(result, res_name);
fun.to_msl()?
} else if context.expression.resolve_type(value).scalar_width() == Some(8) {
fun.to_msl_64_bit()?
} else {
fun.to_msl()?
};
self.put_atomic_operation(pointer, fun_str, value, &context.expression)?; self.put_atomic_operation(pointer, fun_str, value, &context.expression)?;
// done // done
writeln!(self.out, ";")?; writeln!(self.out, ";")?;
@ -5914,3 +5925,31 @@ fn test_stack_size() {
} }
} }
} }
impl crate::AtomicFunction {
fn to_msl(self) -> Result<&'static str, Error> {
Ok(match self {
Self::Add => "fetch_add",
Self::Subtract => "fetch_sub",
Self::And => "fetch_and",
Self::InclusiveOr => "fetch_or",
Self::ExclusiveOr => "fetch_xor",
Self::Min => "fetch_min",
Self::Max => "fetch_max",
Self::Exchange { compare: None } => "exchange",
Self::Exchange { compare: Some(_) } => Err(Error::FeatureNotImplemented(
"atomic CompareExchange".to_string(),
))?,
})
}
fn to_msl_64_bit(self) -> Result<&'static str, Error> {
Ok(match self {
Self::Min => "min",
Self::Max => "max",
_ => Err(Error::FeatureNotImplemented(
"64-bit atomic operation other than min/max".to_string(),
))?,
})
}
}

4
naga/src/back/pipeline_constants.rs
View File

@ -617,7 +617,9 @@ fn adjust_stmt(new_pos: &[Handle<Expression>], stmt: &mut Statement) {
} => { } => {
adjust(pointer); adjust(pointer);
adjust(value); adjust(value);
adjust(result); if let Some(ref mut result) = *result {
adjust(result);
}
match *fun { match *fun {
crate::AtomicFunction::Exchange { crate::AtomicFunction::Exchange {
compare: Some(ref mut compare), compare: Some(ref mut compare),

10
naga/src/back/spv/block.rs
View File

@ -2423,9 +2423,15 @@ impl<'w> BlockContext<'w> {
result, result,
} => { } => {
let id = self.gen_id(); let id = self.gen_id();
let result_type_id = self.get_expression_type_id(&self.fun_info[result].ty); // Compare-and-exchange operations produce a struct result,
// so use `result`'s type if it is available. For no-result
// operations, fall back to `value`'s type.
let result_type_id =
self.get_expression_type_id(&self.fun_info[result.unwrap_or(value)].ty);
self.cached[result] = id; if let Some(result) = result {
self.cached[result] = id;
}
let pointer_id = let pointer_id =
match self.write_expression_pointer(pointer, &mut block, None)? { match self.write_expression_pointer(pointer, &mut block, None)? {

3
naga/src/back/spv/writer.rs
View File

@ -878,6 +878,9 @@ impl Writer {
crate::TypeInner::RayQuery => { crate::TypeInner::RayQuery => {
self.require_any("Ray Query", &[spirv::Capability::RayQueryKHR])?; self.require_any("Ray Query", &[spirv::Capability::RayQueryKHR])?;
} }
crate::TypeInner::Atomic(crate::Scalar { width: 8, kind: _ }) => {
self.require_any("64 bit integer atomics", &[spirv::Capability::Int64Atomics])?;
}
_ => {} _ => {}
} }
Ok(()) Ok(())

8
naga/src/back/wgsl/writer.rs
View File

@ -754,9 +754,11 @@ impl<W: Write> Writer<W> {
result, result,
} => { } => {
write!(self.out, "{level}")?; write!(self.out, "{level}")?;
let res_name = format!("{}{}", back::BAKE_PREFIX, result.index()); if let Some(result) = result {
self.start_named_expr(module, result, func_ctx, &res_name)?; let res_name = format!("{}{}", back::BAKE_PREFIX, result.index());
self.named_expressions.insert(result, res_name); self.start_named_expr(module, result, func_ctx, &res_name)?;
self.named_expressions.insert(result, res_name);
}
let fun_str = fun.to_wgsl(); let fun_str = fun.to_wgsl();
write!(self.out, "atomic{fun_str}(")?; write!(self.out, "atomic{fun_str}(")?;

8
naga/src/compact/statements.rs
View File

@ -75,7 +75,9 @@ impl FunctionTracer<'_> {
self.expressions_used.insert(pointer); self.expressions_used.insert(pointer);
self.trace_atomic_function(fun); self.trace_atomic_function(fun);
self.expressions_used.insert(value); self.expressions_used.insert(value);
self.expressions_used.insert(result); if let Some(result) = result {
self.expressions_used.insert(result);
}
} }
St::WorkGroupUniformLoad { pointer, result } => { St::WorkGroupUniformLoad { pointer, result } => {
self.expressions_used.insert(pointer); self.expressions_used.insert(pointer);
@ -255,7 +257,9 @@ impl FunctionMap {
adjust(pointer); adjust(pointer);
self.adjust_atomic_function(fun); self.adjust_atomic_function(fun);
adjust(value); adjust(value);
adjust(result); if let Some(ref mut result) = *result {
adjust(result);
}
} }
St::WorkGroupUniformLoad { St::WorkGroupUniformLoad {
ref mut pointer, ref mut pointer,

3
naga/src/front/spv/mod.rs
View File

@ -63,6 +63,7 @@ pub const SUPPORTED_CAPABILITIES: &[spirv::Capability] = &[
spirv::Capability::Int8, spirv::Capability::Int8,
spirv::Capability::Int16, spirv::Capability::Int16,
spirv::Capability::Int64, spirv::Capability::Int64,
spirv::Capability::Int64Atomics,
spirv::Capability::Float16, spirv::Capability::Float16,
spirv::Capability::Float64, spirv::Capability::Float64,
spirv::Capability::Geometry, spirv::Capability::Geometry,
@ -4028,7 +4029,7 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
pointer: p_lexp_handle, pointer: p_lexp_handle,
fun: crate::AtomicFunction::Add, fun: crate::AtomicFunction::Add,
value: one_lexp_handle, value: one_lexp_handle,
result: r_lexp_handle, result: Some(r_lexp_handle),
}; };
block.push(stmt, span); block.push(stmt, span);
} }

4
naga/src/front/type_gen.rs
View File

@ -291,10 +291,10 @@ impl crate::Module {
name: Some("exchanged".to_string()), name: Some("exchanged".to_string()),
ty: bool_ty, ty: bool_ty,
binding: None, binding: None,
offset: 4, offset: scalar.width as u32,
}, },
], ],
span: 8, span: scalar.width as u32 * 2,
}, },
} }
} }

47
naga/src/front/wgsl/lower/mod.rs
View File

@ -1491,6 +1491,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
function, function,
arguments, arguments,
&mut ctx.as_expression(block, &mut emitter), &mut ctx.as_expression(block, &mut emitter),
true,
)?; )?;
block.extend(emitter.finish(&ctx.function.expressions)); block.extend(emitter.finish(&ctx.function.expressions));
return Ok(()); return Ok(());
@ -1747,7 +1748,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
ref arguments, ref arguments,
} => { } => {
let handle = self let handle = self
.call(span, function, arguments, ctx)? .call(span, function, arguments, ctx, false)?
.ok_or(Error::FunctionReturnsVoid(function.span))?; .ok_or(Error::FunctionReturnsVoid(function.span))?;
return Ok(Typed::Plain(handle)); return Ok(Typed::Plain(handle));
} }
@ -1941,6 +1942,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
function: &ast::Ident<'source>, function: &ast::Ident<'source>,
arguments: &[Handle<ast::Expression<'source>>], arguments: &[Handle<ast::Expression<'source>>],
ctx: &mut ExpressionContext<'source, '_, '_>, ctx: &mut ExpressionContext<'source, '_, '_>,
is_statement: bool,
) -> Result<Option<Handle<crate::Expression>>, Error<'source>> { ) -> Result<Option<Handle<crate::Expression>>, Error<'source>> {
match ctx.globals.get(function.name) { match ctx.globals.get(function.name) {
Some(&LoweredGlobalDecl::Type(ty)) => { Some(&LoweredGlobalDecl::Type(ty)) => {
@ -2086,7 +2088,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
self.subgroup_gather_helper(span, mode, arguments, ctx)?, self.subgroup_gather_helper(span, mode, arguments, ctx)?,
)); ));
} else if let Some(fun) = crate::AtomicFunction::map(function.name) { } else if let Some(fun) = crate::AtomicFunction::map(function.name) {
return Ok(Some(self.atomic_helper(span, fun, arguments, ctx)?)); return self.atomic_helper(span, fun, arguments, is_statement, ctx);
} else { } else {
match function.name { match function.name {
"select" => { "select" => {
@ -2168,7 +2170,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
compare: Some(compare), compare: Some(compare),
}, },
value, value,
result, result: Some(result),
}, },
span, span,
); );
@ -2459,25 +2461,38 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
span: Span, span: Span,
fun: crate::AtomicFunction, fun: crate::AtomicFunction,
args: &[Handle<ast::Expression<'source>>], args: &[Handle<ast::Expression<'source>>],
is_statement: bool,
ctx: &mut ExpressionContext<'source, '_, '_>, ctx: &mut ExpressionContext<'source, '_, '_>,
) -> Result<Handle<crate::Expression>, Error<'source>> { ) -> Result<Option<Handle<crate::Expression>>, Error<'source>> {
let mut args = ctx.prepare_args(args, 2, span); let mut args = ctx.prepare_args(args, 2, span);
let pointer = self.atomic_pointer(args.next()?, ctx)?; let pointer = self.atomic_pointer(args.next()?, ctx)?;
let value = self.expression(args.next()?, ctx)?;
let value = args.next()?; let value_inner = resolve_inner!(ctx, value);
let value = self.expression(value, ctx)?;
let ty = ctx.register_type(value)?;
args.finish()?; args.finish()?;
let result = ctx.interrupt_emitter( // If we don't use the return value of a 64-bit `min` or `max`
crate::Expression::AtomicResult { // operation, generate a no-result form of the `Atomic` statement, so
ty, // that we can pass validation with only `SHADER_INT64_ATOMIC_MIN_MAX`
comparison: false, // whenever possible.
}, let is_64_bit_min_max =
span, matches!(fun, crate::AtomicFunction::Min | crate::AtomicFunction::Max)
)?; && matches!(
*value_inner,
crate::TypeInner::Scalar(crate::Scalar { width: 8, .. })
);
let result = if is_64_bit_min_max && is_statement {
None
} else {
let ty = ctx.register_type(value)?;
Some(ctx.interrupt_emitter(
crate::Expression::AtomicResult {
ty,
comparison: false,
},
span,
)?)
};
let rctx = ctx.runtime_expression_ctx(span)?; let rctx = ctx.runtime_expression_ctx(span)?;
rctx.block.push( rctx.block.push(
crate::Statement::Atomic { crate::Statement::Atomic {

76
naga/src/lib.rs
View File

@ -1612,8 +1612,29 @@ pub enum Expression {
}, },
/// Result of calling another function. /// Result of calling another function.
CallResult(Handle<Function>), CallResult(Handle<Function>),
/// Result of an atomic operation. /// Result of an atomic operation.
///
/// This expression must be referred to by the [`result`] field of exactly one
/// [`Atomic`][stmt] statement somewhere in the same function. Let `T` be the
/// scalar type contained by the [`Atomic`][type] value that the statement
/// operates on.
///
/// If `comparison` is `false`, then `ty` must be the scalar type `T`.
///
/// If `comparison` is `true`, then `ty` must be a [`Struct`] with two members:
///
/// - A member named `old_value`, whose type is `T`, and
///
/// - A member named `exchanged`, of type [`BOOL`].
///
/// [`result`]: Statement::Atomic::result
/// [stmt]: Statement::Atomic
/// [type]: TypeInner::Atomic
/// [`Struct`]: TypeInner::Struct
/// [`BOOL`]: Scalar::BOOL
AtomicResult { ty: Handle<Type>, comparison: bool }, AtomicResult { ty: Handle<Type>, comparison: bool },
/// Result of a [`WorkGroupUniformLoad`] statement. /// Result of a [`WorkGroupUniformLoad`] statement.
/// ///
/// [`WorkGroupUniformLoad`]: Statement::WorkGroupUniformLoad /// [`WorkGroupUniformLoad`]: Statement::WorkGroupUniformLoad
@ -1890,15 +1911,66 @@ pub enum Statement {
/// Atomic function. /// Atomic function.
Atomic { Atomic {
/// Pointer to an atomic value. /// Pointer to an atomic value.
///
/// This must be a [`Pointer`] to an [`Atomic`] value. The atomic's
/// scalar type may be [`I32`] or [`U32`].
///
/// If [`SHADER_INT64_ATOMIC_MIN_MAX`] or [`SHADER_INT64_ATOMIC_ALL_OPS`] are
/// enabled, this may also be [`I64`] or [`U64`].
///
/// [`Pointer`]: TypeInner::Pointer
/// [`Atomic`]: TypeInner::Atomic
/// [`I32`]: Scalar::I32
/// [`U32`]: Scalar::U32
/// [`SHADER_INT64_ATOMIC_MIN_MAX`]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_MIN_MAX
/// [`SHADER_INT64_ATOMIC_ALL_OPS`]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_ALL_OPS
/// [`I64`]: Scalar::I64
/// [`U64`]: Scalar::U64
pointer: Handle<Expression>, pointer: Handle<Expression>,
/// Function to run on the atomic.
/// Function to run on the atomic value.
///
/// If [`pointer`] refers to a 64-bit atomic value, then:
///
/// - The [`SHADER_INT64_ATOMIC_ALL_OPS`] capability allows any [`AtomicFunction`]
/// value here.
///
/// - The [`SHADER_INT64_ATOMIC_MIN_MAX`] capability allows
/// [`AtomicFunction::Min`] and [`AtomicFunction::Max`] here.
///
/// - If neither of those capabilities are present, then 64-bit scalar
/// atomics are not allowed.
///
/// [`pointer`]: Statement::Atomic::pointer
/// [`SHADER_INT64_ATOMIC_MIN_MAX`]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_MIN_MAX
/// [`SHADER_INT64_ATOMIC_ALL_OPS`]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_ALL_OPS
fun: AtomicFunction, fun: AtomicFunction,
/// Value to use in the function. /// Value to use in the function.
///
/// This must be a scalar of the same type as [`pointer`]'s atomic's scalar type.
///
/// [`pointer`]: Statement::Atomic::pointer
value: Handle<Expression>, value: Handle<Expression>,
/// [`AtomicResult`] expression representing this function's result. /// [`AtomicResult`] expression representing this function's result.
/// ///
/// If [`fun`] is [`Exchange { compare: None }`], this must be `Some`,
/// as otherwise that operation would be equivalent to a simple [`Store`]
/// to the atomic.
///
/// Otherwise, this may be `None` if the return value of the operation is not needed.
///
/// If `pointer` refers to a 64-bit atomic value, [`SHADER_INT64_ATOMIC_MIN_MAX`]
/// is enabled, and [`SHADER_INT64_ATOMIC_ALL_OPS`] is not, this must be `None`.
///
/// [`AtomicResult`]: crate::Expression::AtomicResult /// [`AtomicResult`]: crate::Expression::AtomicResult
result: Handle<Expression>, /// [`fun`]: Statement::Atomic::fun
/// [`Store`]: Statement::Store
/// [`Exchange { compare: None }`]: AtomicFunction::Exchange
/// [`SHADER_INT64_ATOMIC_MIN_MAX`]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_MIN_MAX
/// [`SHADER_INT64_ATOMIC_ALL_OPS`]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_ALL_OPS
result: Option<Handle<Expression>>,
}, },
/// Load uniformly from a uniform pointer in the workgroup address space. /// Load uniformly from a uniform pointer in the workgroup address space.
/// ///

36
naga/src/valid/expression.rs
View File

@ -1,7 +1,4 @@
use super::{ use super::{compose::validate_compose, FunctionInfo, ModuleInfo, ShaderStages, TypeFlags};
compose::validate_compose, validate_atomic_compare_exchange_struct, FunctionInfo, ModuleInfo,
ShaderStages, TypeFlags,
};
use crate::arena::UniqueArena; use crate::arena::UniqueArena;
use crate::{ use crate::{
@ -114,8 +111,6 @@ pub enum ExpressionError {
WrongArgumentCount(crate::MathFunction), WrongArgumentCount(crate::MathFunction),
#[error("Argument [{1}] to {0:?} as expression {2:?} has an invalid type.")] #[error("Argument [{1}] to {0:?} as expression {2:?} has an invalid type.")]
InvalidArgumentType(crate::MathFunction, u32, Handle<crate::Expression>), InvalidArgumentType(crate::MathFunction, u32, Handle<crate::Expression>),
#[error("Atomic result type can't be {0:?}")]
InvalidAtomicResultType(Handle<crate::Type>),
#[error( #[error(
"workgroupUniformLoad result type can't be {0:?}. It can only be a constructible type." "workgroupUniformLoad result type can't be {0:?}. It can only be a constructible type."
)] )]
@ -1582,30 +1577,11 @@ impl super::Validator {
ShaderStages::all() ShaderStages::all()
} }
E::CallResult(function) => mod_info.functions[function.index()].available_stages, E::CallResult(function) => mod_info.functions[function.index()].available_stages,
E::AtomicResult { ty, comparison } => { E::AtomicResult { .. } => {
let scalar_predicate = |ty: &crate::TypeInner| match ty { // These expressions are validated when we check the `Atomic` statement
&crate::TypeInner::Scalar( // that refers to them, because we have all the information we need at
scalar @ Sc { // that point. The checks driven by `Validator::needs_visit` ensure
kind: crate::ScalarKind::Uint | crate::ScalarKind::Sint, // that this expression is indeed visited by one `Atomic` statement.
..
},
) => self.check_width(scalar).is_ok(),
_ => false,
};
let good = match &module.types[ty].inner {
ty if !comparison => scalar_predicate(ty),
&crate::TypeInner::Struct { ref members, .. } if comparison => {
validate_atomic_compare_exchange_struct(
&module.types,
members,
scalar_predicate,
)
}
_ => false,
};
if !good {
return Err(ExpressionError::InvalidAtomicResultType(ty));
}
ShaderStages::all() ShaderStages::all()
} }
E::WorkGroupUniformLoadResult { ty } => { E::WorkGroupUniformLoadResult { ty } => {

230
naga/src/valid/function.rs
View File

@ -43,10 +43,22 @@ pub enum CallError {
pub enum AtomicError { pub enum AtomicError {
#[error("Pointer {0:?} to atomic is invalid.")] #[error("Pointer {0:?} to atomic is invalid.")]
InvalidPointer(Handle<crate::Expression>), InvalidPointer(Handle<crate::Expression>),
#[error("Address space {0:?} does not support 64bit atomics.")]
InvalidAddressSpace(crate::AddressSpace),
#[error("Operand {0:?} has invalid type.")] #[error("Operand {0:?} has invalid type.")]
InvalidOperand(Handle<crate::Expression>), InvalidOperand(Handle<crate::Expression>),
#[error("Result expression {0:?} is not an `AtomicResult` expression")]
InvalidResultExpression(Handle<crate::Expression>),
#[error("Result expression {0:?} is marked as an `exchange`")]
ResultExpressionExchange(Handle<crate::Expression>),
#[error("Result expression {0:?} is not marked as an `exchange`")]
ResultExpressionNotExchange(Handle<crate::Expression>),
#[error("Result type for {0:?} doesn't match the statement")] #[error("Result type for {0:?} doesn't match the statement")]
ResultTypeMismatch(Handle<crate::Expression>), ResultTypeMismatch(Handle<crate::Expression>),
#[error("Exchange operations must return a value")]
MissingReturnValue,
#[error("Capability {0:?} is required")]
MissingCapability(super::Capabilities),
#[error("Result expression {0:?} is populated by multiple `Atomic` statements")] #[error("Result expression {0:?} is populated by multiple `Atomic` statements")]
ResultAlreadyPopulated(Handle<crate::Expression>), ResultAlreadyPopulated(Handle<crate::Expression>),
} }
@ -350,79 +362,189 @@ impl super::Validator {
pointer: Handle<crate::Expression>, pointer: Handle<crate::Expression>,
fun: &crate::AtomicFunction, fun: &crate::AtomicFunction,
value: Handle<crate::Expression>, value: Handle<crate::Expression>,
result: Handle<crate::Expression>, result: Option<Handle<crate::Expression>>,
span: crate::Span,
context: &BlockContext, context: &BlockContext,
) -> Result<(), WithSpan<FunctionError>> { ) -> Result<(), WithSpan<FunctionError>> {
// The `pointer` operand must be a pointer to an atomic value.
let pointer_inner = context.resolve_type(pointer, &self.valid_expression_set)?; let pointer_inner = context.resolve_type(pointer, &self.valid_expression_set)?;
let ptr_scalar = match *pointer_inner { let crate::TypeInner::Pointer {
crate::TypeInner::Pointer { base, .. } => match context.types[base].inner { base: pointer_base,
crate::TypeInner::Atomic(scalar) => scalar, space: pointer_space,
ref other => { } = *pointer_inner
log::error!("Atomic pointer to type {:?}", other); else {
return Err(AtomicError::InvalidPointer(pointer) log::error!("Atomic operation on type {:?}", *pointer_inner);
.with_span_handle(pointer, context.expressions) return Err(AtomicError::InvalidPointer(pointer)
.into_other()); .with_span_handle(pointer, context.expressions)
} .into_other());
}, };
ref other => { let crate::TypeInner::Atomic(pointer_scalar) = context.types[pointer_base].inner else {
log::error!("Atomic on type {:?}", other); log::error!(
return Err(AtomicError::InvalidPointer(pointer) "Atomic pointer to type {:?}",
.with_span_handle(pointer, context.expressions) context.types[pointer_base].inner
.into_other()); );
} return Err(AtomicError::InvalidPointer(pointer)
.with_span_handle(pointer, context.expressions)
.into_other());
}; };
// The `value` operand must be a scalar of the same type as the atomic.
let value_inner = context.resolve_type(value, &self.valid_expression_set)?; let value_inner = context.resolve_type(value, &self.valid_expression_set)?;
match *value_inner { let crate::TypeInner::Scalar(value_scalar) = *value_inner else {
crate::TypeInner::Scalar(scalar) if scalar == ptr_scalar => {} log::error!("Atomic operand type {:?}", *value_inner);
ref other => { return Err(AtomicError::InvalidOperand(value)
log::error!("Atomic operand type {:?}", other); .with_span_handle(value, context.expressions)
return Err(AtomicError::InvalidOperand(value) .into_other());
};
if pointer_scalar != value_scalar {
log::error!("Atomic operand type {:?}", *value_inner);
return Err(AtomicError::InvalidOperand(value)
.with_span_handle(value, context.expressions)
.into_other());
}
// Check for the special restrictions on 64-bit atomic operations.
//
// We don't need to consider other widths here: this function has already checked
// that `pointer`'s type is an `Atomic`, and `validate_type` has already checked
// that that `Atomic` type has a permitted scalar width.
if pointer_scalar.width == 8 {
// `Capabilities::SHADER_INT64_ATOMIC_ALL_OPS` enables all sorts of 64-bit
// atomic operations.
if self
.capabilities
.contains(super::Capabilities::SHADER_INT64_ATOMIC_ALL_OPS)
{
// okay
} else {
// `Capabilities::SHADER_INT64_ATOMIC_MIN_MAX` allows `Min` and
// `Max` on operations in `Storage`, without a return value.
if matches!(
*fun,
crate::AtomicFunction::Min | crate::AtomicFunction::Max
) && matches!(pointer_space, crate::AddressSpace::Storage { .. })
&& result.is_none()
{
if !self
.capabilities
.contains(super::Capabilities::SHADER_INT64_ATOMIC_MIN_MAX)
{
log::error!("Int64 min-max atomic operations are not supported");
return Err(AtomicError::MissingCapability(
super::Capabilities::SHADER_INT64_ATOMIC_MIN_MAX,
)
.with_span_handle(value, context.expressions)
.into_other());
}
} else {
// Otherwise, we require the full 64-bit atomic capability.
log::error!("Int64 atomic operations are not supported");
return Err(AtomicError::MissingCapability(
super::Capabilities::SHADER_INT64_ATOMIC_ALL_OPS,
)
.with_span_handle(value, context.expressions) .with_span_handle(value, context.expressions)
.into_other()); .into_other());
}
} }
} }
if let crate::AtomicFunction::Exchange { compare: Some(cmp) } = *fun { // The result expression must be appropriate to the operation.
if context.resolve_type(cmp, &self.valid_expression_set)? != value_inner { match result {
log::error!("Atomic exchange comparison has a different type from the value"); Some(result) => {
return Err(AtomicError::InvalidOperand(cmp) // The `result` handle must refer to an `AtomicResult` expression.
.with_span_handle(cmp, context.expressions) let crate::Expression::AtomicResult {
.into_other()); ty: result_ty,
} comparison,
} } = context.expressions[result]
else {
return Err(AtomicError::InvalidResultExpression(result)
.with_span_handle(result, context.expressions)
.into_other());
};
self.emit_expression(result, context)?; // Note that this expression has been visited by the proper kind
match context.expressions[result] { // of statement.
crate::Expression::AtomicResult { ty, comparison }
if {
let scalar_predicate =
|ty: &crate::TypeInner| *ty == crate::TypeInner::Scalar(ptr_scalar);
match &context.types[ty].inner {
ty if !comparison => scalar_predicate(ty),
&crate::TypeInner::Struct { ref members, .. } if comparison => {
validate_atomic_compare_exchange_struct(
context.types,
members,
scalar_predicate,
)
}
_ => false,
}
} =>
{
if !self.needs_visit.remove(result.index()) { if !self.needs_visit.remove(result.index()) {
return Err(AtomicError::ResultAlreadyPopulated(result) return Err(AtomicError::ResultAlreadyPopulated(result)
.with_span_handle(result, context.expressions) .with_span_handle(result, context.expressions)
.into_other()); .into_other());
} }
// The constraints on the result type depend on the atomic function.
if let crate::AtomicFunction::Exchange {
compare: Some(compare),
} = *fun
{
// The comparison value must be a scalar of the same type as the
// atomic we're operating on.
let compare_inner =
context.resolve_type(compare, &self.valid_expression_set)?;
if !compare_inner.equivalent(value_inner, context.types) {
log::error!(
"Atomic exchange comparison has a different type from the value"
);
return Err(AtomicError::InvalidOperand(compare)
.with_span_handle(compare, context.expressions)
.into_other());
}
// The result expression must be an `__atomic_compare_exchange_result`
// struct whose `old_value` member is of the same type as the atomic
// we're operating on.
let crate::TypeInner::Struct { ref members, .. } =
context.types[result_ty].inner
else {
return Err(AtomicError::ResultTypeMismatch(result)
.with_span_handle(result, context.expressions)
.into_other());
};
if !validate_atomic_compare_exchange_struct(
context.types,
members,
|ty: &crate::TypeInner| *ty == crate::TypeInner::Scalar(pointer_scalar),
) {
return Err(AtomicError::ResultTypeMismatch(result)
.with_span_handle(result, context.expressions)
.into_other());
}
// The result expression must be for a comparison operation.
if !comparison {
return Err(AtomicError::ResultExpressionNotExchange(result)
.with_span_handle(result, context.expressions)
.into_other());
}
} else {
// The result expression must be a scalar of the same type as the
// atomic we're operating on.
let result_inner = &context.types[result_ty].inner;
if !result_inner.equivalent(value_inner, context.types) {
return Err(AtomicError::ResultTypeMismatch(result)
.with_span_handle(result, context.expressions)
.into_other());
}
// The result expression must not be for a comparison.
if comparison {
return Err(AtomicError::ResultExpressionExchange(result)
.with_span_handle(result, context.expressions)
.into_other());
}
}
self.emit_expression(result, context)?;
} }
_ => {
return Err(AtomicError::ResultTypeMismatch(result) None => {
.with_span_handle(result, context.expressions) // Exchange operations must always produce a value.
.into_other()) if let crate::AtomicFunction::Exchange { compare: None } = *fun {
log::error!("Atomic exchange's value is unused");
return Err(AtomicError::MissingReturnValue
.with_span_static(span, "atomic exchange operation")
.into_other());
}
} }
} }
Ok(()) Ok(())
} }
fn validate_subgroup_operation( fn validate_subgroup_operation(
@ -1017,7 +1139,7 @@ impl super::Validator {
value, value,
result, result,
} => { } => {
self.validate_atomic(pointer, fun, value, result, context)?; self.validate_atomic(pointer, fun, value, result, span, context)?;
} }
S::WorkGroupUniformLoad { pointer, result } => { S::WorkGroupUniformLoad { pointer, result } => {
stages &= super::ShaderStages::COMPUTE; stages &= super::ShaderStages::COMPUTE;

4
naga/src/valid/handles.rs
View File

@ -530,7 +530,9 @@ impl super::Validator {
crate::AtomicFunction::Exchange { compare } => validate_expr_opt(compare)?, crate::AtomicFunction::Exchange { compare } => validate_expr_opt(compare)?,
}; };
validate_expr(value)?; validate_expr(value)?;
validate_expr(result)?; if let Some(result) = result {
validate_expr(result)?;
}
Ok(()) Ok(())
} }
crate::Statement::WorkGroupUniformLoad { pointer, result } => { crate::Statement::WorkGroupUniformLoad { pointer, result } => {

12
naga/src/valid/mod.rs
View File

@ -127,6 +127,18 @@ bitflags::bitflags! {
const SUBGROUP = 0x10000; const SUBGROUP = 0x10000;
/// Support for subgroup barriers. /// Support for subgroup barriers.
const SUBGROUP_BARRIER = 0x20000; const SUBGROUP_BARRIER = 0x20000;
/// Support for [`AtomicFunction::Min`] and [`AtomicFunction::Max`] on
/// 64-bit integers in the [`Storage`] address space, when the return
/// value is not used.
///
/// This is the only 64-bit atomic functionality available on Metal 3.1.
///
/// [`AtomicFunction::Min`]: crate::AtomicFunction::Min
/// [`AtomicFunction::Max`]: crate::AtomicFunction::Max
/// [`Storage`]: crate::AddressSpace::Storage
const SHADER_INT64_ATOMIC_MIN_MAX = 0x40000;
/// Support for all atomic operations on 64-bit integers.
const SHADER_INT64_ATOMIC_ALL_OPS = 0x80000;
} }
} }

24
naga/src/valid/type.rs
View File

@ -360,16 +360,28 @@ impl super::Validator {
) )
} }
Ti::Atomic(crate::Scalar { kind, width }) => { Ti::Atomic(crate::Scalar { kind, width }) => {
let good = match kind { match kind {
crate::ScalarKind::Bool crate::ScalarKind::Bool
| crate::ScalarKind::Float | crate::ScalarKind::Float
| crate::ScalarKind::AbstractInt | crate::ScalarKind::AbstractInt
| crate::ScalarKind::AbstractFloat => false, | crate::ScalarKind::AbstractFloat => {
crate::ScalarKind::Sint | crate::ScalarKind::Uint => width == 4, return Err(TypeError::InvalidAtomicWidth(kind, width))
}
crate::ScalarKind::Sint | crate::ScalarKind::Uint => {
if width == 8 {
if !self.capabilities.intersects(
Capabilities::SHADER_INT64_ATOMIC_ALL_OPS
| Capabilities::SHADER_INT64_ATOMIC_MIN_MAX,
) {
return Err(TypeError::MissingCapability(
Capabilities::SHADER_INT64_ATOMIC_ALL_OPS,
));
}
} else if width != 4 {
return Err(TypeError::InvalidAtomicWidth(kind, width));
}
}
}; };
if !good {
return Err(TypeError::InvalidAtomicWidth(kind, width));
}
TypeInfo::new( TypeInfo::new(
TypeFlags::DATA | TypeFlags::SIZED | TypeFlags::HOST_SHAREABLE, TypeFlags::DATA | TypeFlags::SIZED | TypeFlags::HOST_SHAREABLE,
Alignment::from_width(width), Alignment::from_width(width),

15
naga/tests/in/atomicCompareExchange-int64.param.ron Normal file
View File

@ -0,0 +1,15 @@
(
god_mode: true,
spv: (
version: (1, 0),
capabilities: [ Int64, Int64Atomics ],
),
hlsl: (
shader_model: V6_6,
binding_map: {},
fake_missing_bindings: true,
special_constants_binding: Some((space: 1, register: 0)),
push_constants_target: Some((space: 0, register: 0)),
zero_initialize_workgroup_memory: true,
),
)

34
naga/tests/in/atomicCompareExchange-int64.wgsl Normal file
View File

@ -0,0 +1,34 @@
const SIZE: u32 = 128u;
@group(0) @binding(0)
var<storage,read_write> arr_i64: array<atomic<i64>, SIZE>;
@group(0) @binding(1)
var<storage,read_write> arr_u64: array<atomic<u64>, SIZE>;
@compute @workgroup_size(1)
fn test_atomic_compare_exchange_i64() {
for(var i = 0u; i < SIZE; i++) {
var old : i64 = atomicLoad(&arr_i64[i]);
var exchanged = false;
while(!exchanged) {
let new_ : i64 = bitcast<i64>(old + 10li);
let result = atomicCompareExchangeWeak(&arr_i64[i], old, new_);
old = result.old_value;
exchanged = result.exchanged;
}
}
}
@compute @workgroup_size(1)
fn test_atomic_compare_exchange_u64() {
for(var i = 0u; i < SIZE; i++) {
var old : u64 = atomicLoad(&arr_u64[i]);
var exchanged = false;
while(!exchanged) {
let new_ : u64 = bitcast<u64>(old + 10lu);
let result = atomicCompareExchangeWeak(&arr_u64[i], old, new_);
old = result.old_value;
exchanged = result.exchanged;
}
}
}

23
naga/tests/in/atomicOps-int64-min-max.param.ron Normal file
View File

@ -0,0 +1,23 @@
(
god_mode: true,
spv: (
version: (1, 0),
capabilities: [ Int64, Int64Atomics ],
),
hlsl: (
shader_model: V6_6,
binding_map: {},
fake_missing_bindings: true,
special_constants_binding: Some((space: 1, register: 0)),
push_constants_target: Some((space: 0, register: 0)),
zero_initialize_workgroup_memory: true,
),
msl: (
lang_version: (2, 4),
per_entry_point_map: {},
inline_samplers: [],
spirv_cross_compatibility: false,
fake_missing_bindings: true,
zero_initialize_workgroup_memory: true,
),
)

27
naga/tests/in/atomicOps-int64-min-max.wgsl Normal file
View File

@ -0,0 +1,27 @@
struct Struct {
atomic_scalar: atomic<u64>,
atomic_arr: array<atomic<u64>, 2>,
}
@group(0) @binding(0)
var<storage, read_write> storage_atomic_scalar: atomic<u64>;
@group(0) @binding(1)
var<storage, read_write> storage_atomic_arr: array<atomic<u64>, 2>;
@group(0) @binding(2)
var<storage, read_write> storage_struct: Struct;
@compute
@workgroup_size(2)
fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
atomicMax(&storage_atomic_scalar, 1lu);
atomicMax(&storage_atomic_arr[1], 1lu);
atomicMax(&storage_struct.atomic_scalar, 1lu);
atomicMax(&storage_struct.atomic_arr[1], 1lu);
workgroupBarrier();
atomicMin(&storage_atomic_scalar, 1lu);
atomicMin(&storage_atomic_arr[1], 1lu);
atomicMin(&storage_struct.atomic_scalar, 1lu);
atomicMin(&storage_struct.atomic_arr[1], 1lu);
}

15
naga/tests/in/atomicOps-int64.param.ron Normal file
View File

@ -0,0 +1,15 @@
(
god_mode: true,
spv: (
version: (1, 0),
capabilities: [ Int64, Int64Atomics ],
),
hlsl: (
shader_model: V6_6,
binding_map: {},
fake_missing_bindings: true,
special_constants_binding: Some((space: 1, register: 0)),
push_constants_target: Some((space: 0, register: 0)),
zero_initialize_workgroup_memory: true,
),
)

141
naga/tests/in/atomicOps-int64.wgsl Normal file
View File

@ -0,0 +1,141 @@
// This test covers the cross product of:
//
// * All int64 atomic operations.
// * On all applicable scopes (storage read-write, workgroup).
// * For all shapes of modeling atomic data.
struct Struct {
atomic_scalar: atomic<u64>,
atomic_arr: array<atomic<i64>, 2>,
}
@group(0) @binding(0)
var<storage, read_write> storage_atomic_scalar: atomic<u64>;
@group(0) @binding(1)
var<storage, read_write> storage_atomic_arr: array<atomic<i64>, 2>;
@group(0) @binding(2)
var<storage, read_write> storage_struct: Struct;
var<workgroup> workgroup_atomic_scalar: atomic<u64>;
var<workgroup> workgroup_atomic_arr: array<atomic<i64>, 2>;
var<workgroup> workgroup_struct: Struct;
@compute
@workgroup_size(2)
fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
atomicStore(&storage_atomic_scalar, 1lu);
atomicStore(&storage_atomic_arr[1], 1li);
atomicStore(&storage_struct.atomic_scalar, 1lu);
atomicStore(&storage_struct.atomic_arr[1], 1li);
atomicStore(&workgroup_atomic_scalar, 1lu);
atomicStore(&workgroup_atomic_arr[1], 1li);
atomicStore(&workgroup_struct.atomic_scalar, 1lu);
atomicStore(&workgroup_struct.atomic_arr[1], 1li);
workgroupBarrier();
let l0 = atomicLoad(&storage_atomic_scalar);
let l1 = atomicLoad(&storage_atomic_arr[1]);
let l2 = atomicLoad(&storage_struct.atomic_scalar);
let l3 = atomicLoad(&storage_struct.atomic_arr[1]);
let l4 = atomicLoad(&workgroup_atomic_scalar);
let l5 = atomicLoad(&workgroup_atomic_arr[1]);
let l6 = atomicLoad(&workgroup_struct.atomic_scalar);
let l7 = atomicLoad(&workgroup_struct.atomic_arr[1]);
workgroupBarrier();
atomicAdd(&storage_atomic_scalar, 1lu);
atomicAdd(&storage_atomic_arr[1], 1li);
atomicAdd(&storage_struct.atomic_scalar, 1lu);
atomicAdd(&storage_struct.atomic_arr[1], 1li);
atomicAdd(&workgroup_atomic_scalar, 1lu);
atomicAdd(&workgroup_atomic_arr[1], 1li);
atomicAdd(&workgroup_struct.atomic_scalar, 1lu);
atomicAdd(&workgroup_struct.atomic_arr[1], 1li);
workgroupBarrier();
atomicSub(&storage_atomic_scalar, 1lu);
atomicSub(&storage_atomic_arr[1], 1li);
atomicSub(&storage_struct.atomic_scalar, 1lu);
atomicSub(&storage_struct.atomic_arr[1], 1li);
atomicSub(&workgroup_atomic_scalar, 1lu);
atomicSub(&workgroup_atomic_arr[1], 1li);
atomicSub(&workgroup_struct.atomic_scalar, 1lu);
atomicSub(&workgroup_struct.atomic_arr[1], 1li);
workgroupBarrier();
atomicMax(&storage_atomic_scalar, 1lu);
atomicMax(&storage_atomic_arr[1], 1li);
atomicMax(&storage_struct.atomic_scalar, 1lu);
atomicMax(&storage_struct.atomic_arr[1], 1li);
atomicMax(&workgroup_atomic_scalar, 1lu);
atomicMax(&workgroup_atomic_arr[1], 1li);
atomicMax(&workgroup_struct.atomic_scalar, 1lu);
atomicMax(&workgroup_struct.atomic_arr[1], 1li);
workgroupBarrier();
atomicMin(&storage_atomic_scalar, 1lu);
atomicMin(&storage_atomic_arr[1], 1li);
atomicMin(&storage_struct.atomic_scalar, 1lu);
atomicMin(&storage_struct.atomic_arr[1], 1li);
atomicMin(&workgroup_atomic_scalar, 1lu);
atomicMin(&workgroup_atomic_arr[1], 1li);
atomicMin(&workgroup_struct.atomic_scalar, 1lu);
atomicMin(&workgroup_struct.atomic_arr[1], 1li);
workgroupBarrier();
atomicAnd(&storage_atomic_scalar, 1lu);
atomicAnd(&storage_atomic_arr[1], 1li);
atomicAnd(&storage_struct.atomic_scalar, 1lu);
atomicAnd(&storage_struct.atomic_arr[1], 1li);
atomicAnd(&workgroup_atomic_scalar, 1lu);
atomicAnd(&workgroup_atomic_arr[1], 1li);
atomicAnd(&workgroup_struct.atomic_scalar, 1lu);
atomicAnd(&workgroup_struct.atomic_arr[1], 1li);
workgroupBarrier();
atomicOr(&storage_atomic_scalar, 1lu);
atomicOr(&storage_atomic_arr[1], 1li);
atomicOr(&storage_struct.atomic_scalar, 1lu);
atomicOr(&storage_struct.atomic_arr[1], 1li);
atomicOr(&workgroup_atomic_scalar, 1lu);
atomicOr(&workgroup_atomic_arr[1], 1li);
atomicOr(&workgroup_struct.atomic_scalar, 1lu);
atomicOr(&workgroup_struct.atomic_arr[1], 1li);
workgroupBarrier();
atomicXor(&storage_atomic_scalar, 1lu);
atomicXor(&storage_atomic_arr[1], 1li);
atomicXor(&storage_struct.atomic_scalar, 1lu);
atomicXor(&storage_struct.atomic_arr[1], 1li);
atomicXor(&workgroup_atomic_scalar, 1lu);
atomicXor(&workgroup_atomic_arr[1], 1li);
atomicXor(&workgroup_struct.atomic_scalar, 1lu);
atomicXor(&workgroup_struct.atomic_arr[1], 1li);
atomicExchange(&storage_atomic_scalar, 1lu);
atomicExchange(&storage_atomic_arr[1], 1li);
atomicExchange(&storage_struct.atomic_scalar, 1lu);
atomicExchange(&storage_struct.atomic_arr[1], 1li);
atomicExchange(&workgroup_atomic_scalar, 1lu);
atomicExchange(&workgroup_atomic_arr[1], 1li);
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);
}

1
naga/tests/in/int64.param.ron
View File

@ -2,6 +2,7 @@
god_mode: true, god_mode: true,
spv: ( spv: (
version: (1, 0), version: (1, 0),
capabilities: [ Int64 ],
), ),
hlsl: ( hlsl: (
shader_model: V6_0, shader_model: V6_0,

30
naga/tests/out/hlsl/atomicOps-int64-min-max.hlsl Normal file
View File

@ -0,0 +1,30 @@
struct NagaConstants {
int first_vertex;
int first_instance;
uint other;
};
ConstantBuffer<NagaConstants> _NagaConstants: register(b0, space1);
struct Struct {
uint64_t atomic_scalar;
uint64_t atomic_arr[2];
};
RWByteAddressBuffer storage_atomic_scalar : register(u0);
RWByteAddressBuffer storage_atomic_arr : register(u1);
RWByteAddressBuffer storage_struct : register(u2);
[numthreads(2, 1, 1)]
void cs_main(uint3 id : SV_GroupThreadID)
{
storage_atomic_scalar.InterlockedMax(0, 1uL);
storage_atomic_arr.InterlockedMax(8, 1uL);
storage_struct.InterlockedMax(0, 1uL);
storage_struct.InterlockedMax(8+8, 1uL);
GroupMemoryBarrierWithGroupSync();
storage_atomic_scalar.InterlockedMin(0, 1uL);
storage_atomic_arr.InterlockedMin(8, 1uL);
storage_struct.InterlockedMin(0, 1uL);
storage_struct.InterlockedMin(8+8, 1uL);
return;
}

12
naga/tests/out/hlsl/atomicOps-int64-min-max.ron Normal file
View File

@ -0,0 +1,12 @@
(
vertex:[
],
fragment:[
],
compute:[
(
entry_point:"cs_main",
target_profile:"cs_6_6",
),
],
)

118
naga/tests/out/hlsl/atomicOps-int64.hlsl Normal file
View File

@ -0,0 +1,118 @@
struct NagaConstants {
int first_vertex;
int first_instance;
uint other;
};
ConstantBuffer<NagaConstants> _NagaConstants: register(b0, space1);
struct Struct {
uint64_t atomic_scalar;
int64_t atomic_arr[2];
};
RWByteAddressBuffer storage_atomic_scalar : register(u0);
RWByteAddressBuffer storage_atomic_arr : register(u1);
RWByteAddressBuffer storage_struct : register(u2);
groupshared uint64_t workgroup_atomic_scalar;
groupshared int64_t 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 = (uint64_t)0;
workgroup_atomic_arr = (int64_t[2])0;
workgroup_struct = (Struct)0;
}
GroupMemoryBarrierWithGroupSync();
storage_atomic_scalar.Store(0, 1uL);
storage_atomic_arr.Store(8, 1L);
storage_struct.Store(0, 1uL);
storage_struct.Store(8+8, 1L);
workgroup_atomic_scalar = 1uL;
workgroup_atomic_arr[1] = 1L;
workgroup_struct.atomic_scalar = 1uL;
workgroup_struct.atomic_arr[1] = 1L;
GroupMemoryBarrierWithGroupSync();
uint64_t l0_ = storage_atomic_scalar.Load<uint64_t>(0);
int64_t l1_ = storage_atomic_arr.Load<int64_t>(8);
uint64_t l2_ = storage_struct.Load<uint64_t>(0);
int64_t l3_ = storage_struct.Load<int64_t>(8+8);
uint64_t l4_ = workgroup_atomic_scalar;
int64_t l5_ = workgroup_atomic_arr[1];
uint64_t l6_ = workgroup_struct.atomic_scalar;
int64_t l7_ = workgroup_struct.atomic_arr[1];
GroupMemoryBarrierWithGroupSync();
uint64_t _e51; storage_atomic_scalar.InterlockedAdd(0, 1uL, _e51);
int64_t _e55; storage_atomic_arr.InterlockedAdd(8, 1L, _e55);
uint64_t _e59; storage_struct.InterlockedAdd(0, 1uL, _e59);
int64_t _e64; storage_struct.InterlockedAdd(8+8, 1L, _e64);
uint64_t _e67; InterlockedAdd(workgroup_atomic_scalar, 1uL, _e67);
int64_t _e71; InterlockedAdd(workgroup_atomic_arr[1], 1L, _e71);
uint64_t _e75; InterlockedAdd(workgroup_struct.atomic_scalar, 1uL, _e75);
int64_t _e80; InterlockedAdd(workgroup_struct.atomic_arr[1], 1L, _e80);
GroupMemoryBarrierWithGroupSync();
uint64_t _e83; storage_atomic_scalar.InterlockedAdd(0, -1uL, _e83);
int64_t _e87; storage_atomic_arr.InterlockedAdd(8, -1L, _e87);
uint64_t _e91; storage_struct.InterlockedAdd(0, -1uL, _e91);
int64_t _e96; storage_struct.InterlockedAdd(8+8, -1L, _e96);
uint64_t _e99; InterlockedAdd(workgroup_atomic_scalar, -1uL, _e99);
int64_t _e103; InterlockedAdd(workgroup_atomic_arr[1], -1L, _e103);
uint64_t _e107; InterlockedAdd(workgroup_struct.atomic_scalar, -1uL, _e107);
int64_t _e112; InterlockedAdd(workgroup_struct.atomic_arr[1], -1L, _e112);
GroupMemoryBarrierWithGroupSync();
storage_atomic_scalar.InterlockedMax(0, 1uL);
storage_atomic_arr.InterlockedMax(8, 1L);
storage_struct.InterlockedMax(0, 1uL);
storage_struct.InterlockedMax(8+8, 1L);
InterlockedMax(workgroup_atomic_scalar, 1uL);
InterlockedMax(workgroup_atomic_arr[1], 1L);
InterlockedMax(workgroup_struct.atomic_scalar, 1uL);
InterlockedMax(workgroup_struct.atomic_arr[1], 1L);
GroupMemoryBarrierWithGroupSync();
storage_atomic_scalar.InterlockedMin(0, 1uL);
storage_atomic_arr.InterlockedMin(8, 1L);
storage_struct.InterlockedMin(0, 1uL);
storage_struct.InterlockedMin(8+8, 1L);
InterlockedMin(workgroup_atomic_scalar, 1uL);
InterlockedMin(workgroup_atomic_arr[1], 1L);
InterlockedMin(workgroup_struct.atomic_scalar, 1uL);
InterlockedMin(workgroup_struct.atomic_arr[1], 1L);
GroupMemoryBarrierWithGroupSync();
uint64_t _e163; storage_atomic_scalar.InterlockedAnd(0, 1uL, _e163);
int64_t _e167; storage_atomic_arr.InterlockedAnd(8, 1L, _e167);
uint64_t _e171; storage_struct.InterlockedAnd(0, 1uL, _e171);
int64_t _e176; storage_struct.InterlockedAnd(8+8, 1L, _e176);
uint64_t _e179; InterlockedAnd(workgroup_atomic_scalar, 1uL, _e179);
int64_t _e183; InterlockedAnd(workgroup_atomic_arr[1], 1L, _e183);
uint64_t _e187; InterlockedAnd(workgroup_struct.atomic_scalar, 1uL, _e187);
int64_t _e192; InterlockedAnd(workgroup_struct.atomic_arr[1], 1L, _e192);
GroupMemoryBarrierWithGroupSync();
uint64_t _e195; storage_atomic_scalar.InterlockedOr(0, 1uL, _e195);
int64_t _e199; storage_atomic_arr.InterlockedOr(8, 1L, _e199);
uint64_t _e203; storage_struct.InterlockedOr(0, 1uL, _e203);
int64_t _e208; storage_struct.InterlockedOr(8+8, 1L, _e208);
uint64_t _e211; InterlockedOr(workgroup_atomic_scalar, 1uL, _e211);
int64_t _e215; InterlockedOr(workgroup_atomic_arr[1], 1L, _e215);
uint64_t _e219; InterlockedOr(workgroup_struct.atomic_scalar, 1uL, _e219);
int64_t _e224; InterlockedOr(workgroup_struct.atomic_arr[1], 1L, _e224);
GroupMemoryBarrierWithGroupSync();
uint64_t _e227; storage_atomic_scalar.InterlockedXor(0, 1uL, _e227);
int64_t _e231; storage_atomic_arr.InterlockedXor(8, 1L, _e231);
uint64_t _e235; storage_struct.InterlockedXor(0, 1uL, _e235);
int64_t _e240; storage_struct.InterlockedXor(8+8, 1L, _e240);
uint64_t _e243; InterlockedXor(workgroup_atomic_scalar, 1uL, _e243);
int64_t _e247; InterlockedXor(workgroup_atomic_arr[1], 1L, _e247);
uint64_t _e251; InterlockedXor(workgroup_struct.atomic_scalar, 1uL, _e251);
int64_t _e256; InterlockedXor(workgroup_struct.atomic_arr[1], 1L, _e256);
uint64_t _e259; storage_atomic_scalar.InterlockedExchange(0, 1uL, _e259);
int64_t _e263; storage_atomic_arr.InterlockedExchange(8, 1L, _e263);
uint64_t _e267; storage_struct.InterlockedExchange(0, 1uL, _e267);
int64_t _e272; storage_struct.InterlockedExchange(8+8, 1L, _e272);
uint64_t _e275; InterlockedExchange(workgroup_atomic_scalar, 1uL, _e275);
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);
return;
}

12
naga/tests/out/hlsl/atomicOps-int64.ron Normal file
View File

@ -0,0 +1,12 @@
(
vertex:[
],
fragment:[
],
compute:[
(
entry_point:"cs_main",
target_profile:"cs_6_6",
),
],
)

2
naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron
View File

@ -116,7 +116,7 @@
compare: Some(3), compare: Some(3),
), ),
value: 4, value: 4,
result: 5, result: Some(5),
), ),
Return( Return(
value: None, value: None,

2
naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron
View File

@ -116,7 +116,7 @@
compare: Some(3), compare: Some(3),
), ),
value: 4, value: 4,
result: 5, result: Some(5),
), ),
Return( Return(
value: None, value: None,

33
naga/tests/out/msl/atomicOps-int64-min-max.msl Normal file
View File

@ -0,0 +1,33 @@
// language: metal2.4
#include <metal_stdlib>
#include <simd/simd.h>
using metal::uint;
struct type_1 {
metal::atomic_ulong inner[2];
};
struct Struct {
metal::atomic_ulong atomic_scalar;
type_1 atomic_arr;
};
struct cs_mainInput {
};
kernel void cs_main(
metal::uint3 id [[thread_position_in_threadgroup]]
, device metal::atomic_ulong& storage_atomic_scalar [[user(fake0)]]
, device type_1& storage_atomic_arr [[user(fake0)]]
, device Struct& storage_struct [[user(fake0)]]
) {
metal::atomic_max_explicit(&storage_atomic_scalar, 1uL, metal::memory_order_relaxed);
metal::atomic_max_explicit(&storage_atomic_arr.inner[1], 1uL, metal::memory_order_relaxed);
metal::atomic_max_explicit(&storage_struct.atomic_scalar, 1uL, metal::memory_order_relaxed);
metal::atomic_max_explicit(&storage_struct.atomic_arr.inner[1], 1uL, metal::memory_order_relaxed);
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
metal::atomic_min_explicit(&storage_atomic_scalar, 1uL, metal::memory_order_relaxed);
metal::atomic_min_explicit(&storage_atomic_arr.inner[1], 1uL, metal::memory_order_relaxed);
metal::atomic_min_explicit(&storage_struct.atomic_scalar, 1uL, metal::memory_order_relaxed);
metal::atomic_min_explicit(&storage_struct.atomic_arr.inner[1], 1uL, metal::memory_order_relaxed);
return;
}

205
naga/tests/out/spv/atomicCompareExchange-int64.spvasm Normal file
View File

@ -0,0 +1,205 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 125
OpCapability Shader
OpCapability Int64Atomics
OpCapability Int64
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %19 "test_atomic_compare_exchange_i64"
OpEntryPoint GLCompute %77 "test_atomic_compare_exchange_u64"
OpExecutionMode %19 LocalSize 1 1 1
OpExecutionMode %77 LocalSize 1 1 1
OpDecorate %5 ArrayStride 8
OpDecorate %8 ArrayStride 8
OpMemberDecorate %10 0 Offset 0
OpMemberDecorate %10 1 Offset 8
OpMemberDecorate %11 0 Offset 0
OpMemberDecorate %11 1 Offset 8
OpDecorate %12 DescriptorSet 0
OpDecorate %12 Binding 0
OpDecorate %13 Block
OpMemberDecorate %13 0 Offset 0
OpDecorate %15 DescriptorSet 0
OpDecorate %15 Binding 1
OpDecorate %16 Block
OpMemberDecorate %16 0 Offset 0
%2 = OpTypeVoid
%3 = OpTypeInt 32 0
%4 = OpTypeInt 64 1
%6 = OpConstant %3 128
%5 = OpTypeArray %4 %6
%7 = OpTypeInt 64 0
%8 = OpTypeArray %7 %6
%9 = OpTypeBool
%10 = OpTypeStruct %4 %9
%11 = OpTypeStruct %7 %9
%13 = OpTypeStruct %5
%14 = OpTypePointer StorageBuffer %13
%12 = OpVariable %14 StorageBuffer
%16 = OpTypeStruct %8
%17 = OpTypePointer StorageBuffer %16
%15 = OpVariable %17 StorageBuffer
%20 = OpTypeFunction %2
%21 = OpTypePointer StorageBuffer %5
%22 = OpConstant %3 0
%24 = OpConstantFalse %9
%25 = OpConstant %4 10
%26 = OpConstant %3 1
%28 = OpTypePointer Function %3
%30 = OpTypePointer Function %4
%31 = OpConstantNull %4
%33 = OpTypePointer Function %9
%34 = OpConstantNull %9
%47 = OpTypePointer StorageBuffer %4
%51 = OpTypeInt 32 1
%50 = OpConstant %51 1
%52 = OpConstant %3 64
%78 = OpTypePointer StorageBuffer %8
%80 = OpConstant %7 10
%83 = OpTypePointer Function %7
%84 = OpConstantNull %7
%86 = OpConstantNull %9
%99 = OpTypePointer StorageBuffer %7
%19 = OpFunction %2 None %20
%18 = OpLabel
%27 = OpVariable %28 Function %22
%29 = OpVariable %30 Function %31
%32 = OpVariable %33 Function %34
%23 = OpAccessChain %21 %12 %22
OpBranch %35
%35 = OpLabel
OpBranch %36
%36 = OpLabel
OpLoopMerge %37 %39 None
OpBranch %38
%38 = OpLabel
%40 = OpLoad %3 %27
%41 = OpULessThan %9 %40 %6
OpSelectionMerge %42 None
OpBranchConditional %41 %42 %43
%43 = OpLabel
OpBranch %37
%42 = OpLabel
OpBranch %44
%44 = OpLabel
%46 = OpLoad %3 %27
%48 = OpAccessChain %47 %23 %46
%49 = OpAtomicLoad %4 %48 %50 %52
OpStore %29 %49
OpStore %32 %24
OpBranch %53
%53 = OpLabel
OpLoopMerge %54 %56 None
OpBranch %55
%55 = OpLabel
%57 = OpLoad %9 %32
%58 = OpLogicalNot %9 %57
OpSelectionMerge %59 None
OpBranchConditional %58 %59 %60
%60 = OpLabel
OpBranch %54
%59 = OpLabel
OpBranch %61
%61 = OpLabel
%63 = OpLoad %4 %29
%64 = OpIAdd %4 %63 %25
%66 = OpLoad %3 %27
%67 = OpLoad %4 %29
%69 = OpAccessChain %47 %23 %66
%70 = OpAtomicCompareExchange %4 %69 %50 %52 %52 %64 %67
%71 = OpIEqual %9 %70 %67
%68 = OpCompositeConstruct %10 %70 %71
%72 = OpCompositeExtract %4 %68 0
OpStore %29 %72
%73 = OpCompositeExtract %9 %68 1
OpStore %32 %73
OpBranch %62
%62 = OpLabel
OpBranch %56
%56 = OpLabel
OpBranch %53
%54 = OpLabel
OpBranch %45
%45 = OpLabel
OpBranch %39
%39 = OpLabel
%74 = OpLoad %3 %27
%75 = OpIAdd %3 %74 %26
OpStore %27 %75
OpBranch %36
%37 = OpLabel
OpReturn
OpFunctionEnd
%77 = OpFunction %2 None %20
%76 = OpLabel
%81 = OpVariable %28 Function %22
%82 = OpVariable %83 Function %84
%85 = OpVariable %33 Function %86
%79 = OpAccessChain %78 %15 %22
OpBranch %87
%87 = OpLabel
OpBranch %88
%88 = OpLabel
OpLoopMerge %89 %91 None
OpBranch %90
%90 = OpLabel
%92 = OpLoad %3 %81
%93 = OpULessThan %9 %92 %6
OpSelectionMerge %94 None
OpBranchConditional %93 %94 %95
%95 = OpLabel
OpBranch %89
%94 = OpLabel
OpBranch %96
%96 = OpLabel
%98 = OpLoad %3 %81
%100 = OpAccessChain %99 %79 %98
%101 = OpAtomicLoad %7 %100 %50 %52
OpStore %82 %101
OpStore %85 %24
OpBranch %102
%102 = OpLabel
OpLoopMerge %103 %105 None
OpBranch %104
%104 = OpLabel
%106 = OpLoad %9 %85
%107 = OpLogicalNot %9 %106
OpSelectionMerge %108 None
OpBranchConditional %107 %108 %109
%109 = OpLabel
OpBranch %103
%108 = OpLabel
OpBranch %110
%110 = OpLabel
%112 = OpLoad %7 %82
%113 = OpIAdd %7 %112 %80
%115 = OpLoad %3 %81
%116 = OpLoad %7 %82
%118 = OpAccessChain %99 %79 %115
%119 = OpAtomicCompareExchange %7 %118 %50 %52 %52 %113 %116
%120 = OpIEqual %9 %119 %116
%117 = OpCompositeConstruct %11 %119 %120
%121 = OpCompositeExtract %7 %117 0
OpStore %82 %121
%122 = OpCompositeExtract %9 %117 1
OpStore %85 %122
OpBranch %111
%111 = OpLabel
OpBranch %105
%105 = OpLabel
OpBranch %102
%103 = OpLabel
OpBranch %97
%97 = OpLabel
OpBranch %91
%91 = OpLabel
%123 = OpLoad %3 %81
%124 = OpIAdd %3 %123 %26
OpStore %81 %124
OpBranch %88
%89 = OpLabel
OpReturn
OpFunctionEnd

82
naga/tests/out/spv/atomicOps-int64-min-max.spvasm Normal file
View File

@ -0,0 +1,82 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 52
OpCapability Shader
OpCapability Int64Atomics
OpCapability Int64
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %22 "cs_main" %19
OpExecutionMode %22 LocalSize 2 1 1
OpDecorate %4 ArrayStride 8
OpMemberDecorate %7 0 Offset 0
OpMemberDecorate %7 1 Offset 8
OpDecorate %9 DescriptorSet 0
OpDecorate %9 Binding 0
OpDecorate %10 Block
OpMemberDecorate %10 0 Offset 0
OpDecorate %12 DescriptorSet 0
OpDecorate %12 Binding 1
OpDecorate %13 Block
OpMemberDecorate %13 0 Offset 0
OpDecorate %15 DescriptorSet 0
OpDecorate %15 Binding 2
OpDecorate %16 Block
OpMemberDecorate %16 0 Offset 0
OpDecorate %19 BuiltIn LocalInvocationId
%2 = OpTypeVoid
%3 = OpTypeInt 64 0
%6 = OpTypeInt 32 0
%5 = OpConstant %6 2
%4 = OpTypeArray %3 %5
%7 = OpTypeStruct %3 %4
%8 = OpTypeVector %6 3
%10 = OpTypeStruct %3
%11 = OpTypePointer StorageBuffer %10
%9 = OpVariable %11 StorageBuffer
%13 = OpTypeStruct %4
%14 = OpTypePointer StorageBuffer %13
%12 = OpVariable %14 StorageBuffer
%16 = OpTypeStruct %7
%17 = OpTypePointer StorageBuffer %16
%15 = OpVariable %17 StorageBuffer
%20 = OpTypePointer Input %8
%19 = OpVariable %20 Input
%23 = OpTypeFunction %2
%24 = OpTypePointer StorageBuffer %3
%25 = OpConstant %6 0
%27 = OpTypePointer StorageBuffer %4
%29 = OpTypePointer StorageBuffer %7
%31 = OpConstant %3 1
%35 = OpTypeInt 32 1
%34 = OpConstant %35 1
%36 = OpConstant %6 64
%38 = OpConstant %6 1
%44 = OpConstant %6 264
%22 = OpFunction %2 None %23
%18 = OpLabel
%21 = OpLoad %8 %19
%26 = OpAccessChain %24 %9 %25
%28 = OpAccessChain %27 %12 %25
%30 = OpAccessChain %29 %15 %25
OpBranch %32
%32 = OpLabel
%33 = OpAtomicUMax %3 %26 %34 %36 %31
%39 = OpAccessChain %24 %28 %38
%37 = OpAtomicUMax %3 %39 %34 %36 %31
%41 = OpAccessChain %24 %30 %25
%40 = OpAtomicUMax %3 %41 %34 %36 %31
%43 = OpAccessChain %24 %30 %38 %38
%42 = OpAtomicUMax %3 %43 %34 %36 %31
OpControlBarrier %5 %5 %44
%45 = OpAtomicUMin %3 %26 %34 %36 %31
%47 = OpAccessChain %24 %28 %38
%46 = OpAtomicUMin %3 %47 %34 %36 %31
%49 = OpAccessChain %24 %30 %25
%48 = OpAtomicUMin %3 %49 %34 %36 %31
%51 = OpAccessChain %24 %30 %38 %38
%50 = OpAtomicUMin %3 %51 %34 %36 %31
OpReturn
OpFunctionEnd

246
naga/tests/out/spv/atomicOps-int64.spvasm Normal file
View File

@ -0,0 +1,246 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 193
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
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
OpDecorate %13 DescriptorSet 0
OpDecorate %13 Binding 1
OpDecorate %14 Block
OpMemberDecorate %14 0 Offset 0
OpDecorate %16 DescriptorSet 0
OpDecorate %16 Binding 2
OpDecorate %17 Block
OpMemberDecorate %17 0 Offset 0
OpDecorate %26 BuiltIn LocalInvocationId
%2 = OpTypeVoid
%3 = OpTypeInt 64 0
%4 = OpTypeInt 64 1
%7 = OpTypeInt 32 0
%6 = OpConstant %7 2
%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
%15 = OpTypePointer StorageBuffer %14
%13 = OpVariable %15 StorageBuffer
%17 = OpTypeStruct %8
%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
OpReturn
OpFunctionEnd

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

@ -0,0 +1,90 @@
const SIZE: u32 = 128u;
@group(0) @binding(0)
var<storage, read_write> arr_i64_: array<atomic<i64>, 128>;
@group(0) @binding(1)
var<storage, read_write> arr_u64_: array<atomic<u64>, 128>;
@compute @workgroup_size(1, 1, 1)
fn test_atomic_compare_exchange_i64_() {
var i: u32 = 0u;
var old: i64;
var exchanged: bool;
loop {
let _e2 = i;
if (_e2 < SIZE) {
} else {
break;
}
{
let _e6 = i;
let _e8 = atomicLoad((&arr_i64_[_e6]));
old = _e8;
exchanged = false;
loop {
let _e12 = exchanged;
if !(_e12) {
} else {
break;
}
{
let _e14 = old;
let new_ = bitcast<i64>((_e14 + 10li));
let _e19 = i;
let _e21 = old;
let _e22 = atomicCompareExchangeWeak((&arr_i64_[_e19]), _e21, new_);
old = _e22.old_value;
exchanged = _e22.exchanged;
}
}
}
continuing {
let _e26 = i;
i = (_e26 + 1u);
}
}
return;
}
@compute @workgroup_size(1, 1, 1)
fn test_atomic_compare_exchange_u64_() {
var i_1: u32 = 0u;
var old_1: u64;
var exchanged_1: bool;
loop {
let _e2 = i_1;
if (_e2 < SIZE) {
} else {
break;
}
{
let _e6 = i_1;
let _e8 = atomicLoad((&arr_u64_[_e6]));
old_1 = _e8;
exchanged_1 = false;
loop {
let _e12 = exchanged_1;
if !(_e12) {
} else {
break;
}
{
let _e14 = old_1;
let new_1 = bitcast<u64>((_e14 + 10lu));
let _e19 = i_1;
let _e21 = old_1;
let _e22 = atomicCompareExchangeWeak((&arr_u64_[_e19]), _e21, new_1);
old_1 = _e22.old_value;
exchanged_1 = _e22.exchanged;
}
}
}
continuing {
let _e26 = i_1;
i_1 = (_e26 + 1u);
}
}
return;
}

25
naga/tests/out/wgsl/atomicOps-int64-min-max.wgsl Normal file
View File

@ -0,0 +1,25 @@
struct Struct {
atomic_scalar: atomic<u64>,
atomic_arr: array<atomic<u64>, 2>,
}
@group(0) @binding(0)
var<storage, read_write> storage_atomic_scalar: atomic<u64>;
@group(0) @binding(1)
var<storage, read_write> storage_atomic_arr: array<atomic<u64>, 2>;
@group(0) @binding(2)
var<storage, read_write> storage_struct: Struct;
@compute @workgroup_size(2, 1, 1)
fn cs_main(@builtin(local_invocation_id) id: vec3<u32>) {
atomicMax((&storage_atomic_scalar), 1lu);
atomicMax((&storage_atomic_arr[1]), 1lu);
atomicMax((&storage_struct.atomic_scalar), 1lu);
atomicMax((&storage_struct.atomic_arr[1]), 1lu);
workgroupBarrier();
atomicMin((&storage_atomic_scalar), 1lu);
atomicMin((&storage_atomic_arr[1]), 1lu);
atomicMin((&storage_struct.atomic_scalar), 1lu);
atomicMin((&storage_struct.atomic_arr[1]), 1lu);
return;
}

107
naga/tests/out/wgsl/atomicOps-int64.wgsl Normal file
View File

@ -0,0 +1,107 @@
struct Struct {
atomic_scalar: atomic<u64>,
atomic_arr: array<atomic<i64>, 2>,
}
@group(0) @binding(0)
var<storage, read_write> storage_atomic_scalar: atomic<u64>;
@group(0) @binding(1)
var<storage, read_write> storage_atomic_arr: array<atomic<i64>, 2>;
@group(0) @binding(2)
var<storage, read_write> storage_struct: Struct;
var<workgroup> workgroup_atomic_scalar: atomic<u64>;
var<workgroup> workgroup_atomic_arr: array<atomic<i64>, 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), 1lu);
atomicStore((&storage_atomic_arr[1]), 1li);
atomicStore((&storage_struct.atomic_scalar), 1lu);
atomicStore((&storage_struct.atomic_arr[1]), 1li);
atomicStore((&workgroup_atomic_scalar), 1lu);
atomicStore((&workgroup_atomic_arr[1]), 1li);
atomicStore((&workgroup_struct.atomic_scalar), 1lu);
atomicStore((&workgroup_struct.atomic_arr[1]), 1li);
workgroupBarrier();
let l0_ = atomicLoad((&storage_atomic_scalar));
let l1_ = atomicLoad((&storage_atomic_arr[1]));
let l2_ = atomicLoad((&storage_struct.atomic_scalar));
let l3_ = atomicLoad((&storage_struct.atomic_arr[1]));
let l4_ = atomicLoad((&workgroup_atomic_scalar));
let l5_ = atomicLoad((&workgroup_atomic_arr[1]));
let l6_ = atomicLoad((&workgroup_struct.atomic_scalar));
let l7_ = atomicLoad((&workgroup_struct.atomic_arr[1]));
workgroupBarrier();
let _e51 = atomicAdd((&storage_atomic_scalar), 1lu);
let _e55 = atomicAdd((&storage_atomic_arr[1]), 1li);
let _e59 = atomicAdd((&storage_struct.atomic_scalar), 1lu);
let _e64 = atomicAdd((&storage_struct.atomic_arr[1]), 1li);
let _e67 = atomicAdd((&workgroup_atomic_scalar), 1lu);
let _e71 = atomicAdd((&workgroup_atomic_arr[1]), 1li);
let _e75 = atomicAdd((&workgroup_struct.atomic_scalar), 1lu);
let _e80 = atomicAdd((&workgroup_struct.atomic_arr[1]), 1li);
workgroupBarrier();
let _e83 = atomicSub((&storage_atomic_scalar), 1lu);
let _e87 = atomicSub((&storage_atomic_arr[1]), 1li);
let _e91 = atomicSub((&storage_struct.atomic_scalar), 1lu);
let _e96 = atomicSub((&storage_struct.atomic_arr[1]), 1li);
let _e99 = atomicSub((&workgroup_atomic_scalar), 1lu);
let _e103 = atomicSub((&workgroup_atomic_arr[1]), 1li);
let _e107 = atomicSub((&workgroup_struct.atomic_scalar), 1lu);
let _e112 = atomicSub((&workgroup_struct.atomic_arr[1]), 1li);
workgroupBarrier();
atomicMax((&storage_atomic_scalar), 1lu);
atomicMax((&storage_atomic_arr[1]), 1li);
atomicMax((&storage_struct.atomic_scalar), 1lu);
atomicMax((&storage_struct.atomic_arr[1]), 1li);
atomicMax((&workgroup_atomic_scalar), 1lu);
atomicMax((&workgroup_atomic_arr[1]), 1li);
atomicMax((&workgroup_struct.atomic_scalar), 1lu);
atomicMax((&workgroup_struct.atomic_arr[1]), 1li);
workgroupBarrier();
atomicMin((&storage_atomic_scalar), 1lu);
atomicMin((&storage_atomic_arr[1]), 1li);
atomicMin((&storage_struct.atomic_scalar), 1lu);
atomicMin((&storage_struct.atomic_arr[1]), 1li);
atomicMin((&workgroup_atomic_scalar), 1lu);
atomicMin((&workgroup_atomic_arr[1]), 1li);
atomicMin((&workgroup_struct.atomic_scalar), 1lu);
atomicMin((&workgroup_struct.atomic_arr[1]), 1li);
workgroupBarrier();
let _e163 = atomicAnd((&storage_atomic_scalar), 1lu);
let _e167 = atomicAnd((&storage_atomic_arr[1]), 1li);
let _e171 = atomicAnd((&storage_struct.atomic_scalar), 1lu);
let _e176 = atomicAnd((&storage_struct.atomic_arr[1]), 1li);
let _e179 = atomicAnd((&workgroup_atomic_scalar), 1lu);
let _e183 = atomicAnd((&workgroup_atomic_arr[1]), 1li);
let _e187 = atomicAnd((&workgroup_struct.atomic_scalar), 1lu);
let _e192 = atomicAnd((&workgroup_struct.atomic_arr[1]), 1li);
workgroupBarrier();
let _e195 = atomicOr((&storage_atomic_scalar), 1lu);
let _e199 = atomicOr((&storage_atomic_arr[1]), 1li);
let _e203 = atomicOr((&storage_struct.atomic_scalar), 1lu);
let _e208 = atomicOr((&storage_struct.atomic_arr[1]), 1li);
let _e211 = atomicOr((&workgroup_atomic_scalar), 1lu);
let _e215 = atomicOr((&workgroup_atomic_arr[1]), 1li);
let _e219 = atomicOr((&workgroup_struct.atomic_scalar), 1lu);
let _e224 = atomicOr((&workgroup_struct.atomic_arr[1]), 1li);
workgroupBarrier();
let _e227 = atomicXor((&storage_atomic_scalar), 1lu);
let _e231 = atomicXor((&storage_atomic_arr[1]), 1li);
let _e235 = atomicXor((&storage_struct.atomic_scalar), 1lu);
let _e240 = atomicXor((&storage_struct.atomic_arr[1]), 1li);
let _e243 = atomicXor((&workgroup_atomic_scalar), 1lu);
let _e247 = atomicXor((&workgroup_atomic_arr[1]), 1li);
let _e251 = atomicXor((&workgroup_struct.atomic_scalar), 1lu);
let _e256 = atomicXor((&workgroup_struct.atomic_arr[1]), 1li);
let _e259 = atomicExchange((&storage_atomic_scalar), 1lu);
let _e263 = atomicExchange((&storage_atomic_arr[1]), 1li);
let _e267 = atomicExchange((&storage_struct.atomic_scalar), 1lu);
let _e272 = atomicExchange((&storage_struct.atomic_arr[1]), 1li);
let _e275 = atomicExchange((&workgroup_atomic_scalar), 1lu);
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);
return;
}

12
naga/tests/snapshots.rs
View File

@ -759,6 +759,18 @@ fn convert_wgsl() {
"padding", "padding",
Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
), ),
(
"atomicOps-int64",
Targets::SPIRV | Targets::HLSL | Targets::WGSL,
),
(
"atomicOps-int64-min-max",
Targets::SPIRV | Targets::METAL | Targets::HLSL | Targets::WGSL,
),
(
"atomicCompareExchange-int64",
Targets::SPIRV | Targets::WGSL,
),
("pointers", Targets::SPIRV | Targets::WGSL), ("pointers", Targets::SPIRV | Targets::WGSL),
( (
"control-flow", "control-flow",

2
naga/tests/validation.rs
View File

@ -75,7 +75,7 @@ fn populate_atomic_result() {
pointer: ex_global, pointer: ex_global,
fun: naga::AtomicFunction::Add, fun: naga::AtomicFunction::Add,
value: ex_42, value: ex_42,
result: ex_result, result: Some(ex_result),
}, },
span, span,
); );

6
tests/tests/shader/mod.rs
View File

@ -155,6 +155,12 @@ impl ShaderTest {
self self
} }
fn output_type(mut self, output_type: String) -> Self {
self.output_type = output_type;
self
}
/// Add another set of possible outputs. If any of the given /// Add another set of possible outputs. If any of the given
/// output values are seen it's considered a success (i.e. this is OR, not AND). /// output values are seen it's considered a success (i.e. this is OR, not AND).
/// ///

99
tests/tests/shader/numeric_builtins.rs
View File

@ -52,6 +52,105 @@ static NUMERIC_BUILTINS: GpuTestConfiguration = GpuTestConfiguration::new()
) )
}); });
fn create_int64_atomic_min_max_test() -> Vec<ShaderTest> {
let mut tests = Vec::new();
let test = ShaderTest::new(
"atomicMax".into(),
"value: u64".into(),
"atomicMin(&output, 0lu); atomicMax(&output, 2lu);".into(),
&[0],
&[2],
)
.output_type("atomic<u64>".into());
tests.push(test);
let test = ShaderTest::new(
"atomicMin".into(),
"value: u64".into(),
"atomicMax(&output, 100lu); atomicMin(&output, 4lu);".into(),
&[0],
&[4],
)
.output_type("atomic<u64>".into());
tests.push(test);
tests
}
#[gpu_test]
static INT64_ATOMIC_MIN_MAX: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.features(wgt::Features::SHADER_INT64 | wgt::Features::SHADER_INT64_ATOMIC_MIN_MAX)
.downlevel_flags(DownlevelFlags::COMPUTE_SHADERS)
.limits(Limits::downlevel_defaults()),
)
.run_async(|ctx| {
shader_input_output_test(
ctx,
InputStorageType::Storage,
create_int64_atomic_min_max_test(),
)
});
fn create_int64_atomic_all_ops_test() -> Vec<ShaderTest> {
let mut tests = Vec::new();
let test = ShaderTest::new(
"atomicAdd".into(),
"value: u64".into(),
"atomicStore(&output, 0lu); atomicAdd(&output, 1lu); atomicAdd(&output, 1lu);".into(),
&[0],
&[2],
)
.output_type("atomic<u64>".into());
tests.push(test);
let test = ShaderTest::new(
"atomicAnd".into(),
"value: u64".into(),
"atomicStore(&output, 31lu); atomicAnd(&output, 30lu); atomicAnd(&output, 3lu);".into(),
&[0],
&[2],
)
.output_type("atomic<u64>".into());
tests.push(test);
let test = ShaderTest::new(
"atomicOr".into(),
"value: u64".into(),
"atomicStore(&output, 0lu); atomicOr(&output, 3lu); atomicOr(&output, 6lu);".into(),
&[0],
&[7],
)
.output_type("atomic<u64>".into());
tests.push(test);
tests
}
#[gpu_test]
static INT64_ATOMIC_ALL_OPS: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.features(wgt::Features::SHADER_INT64 | wgt::Features::SHADER_INT64_ATOMIC_ALL_OPS)
.downlevel_flags(DownlevelFlags::COMPUTE_SHADERS)
.limits(Limits::downlevel_defaults()),
)
.run_async(|ctx| {
shader_input_output_test(
ctx,
InputStorageType::Storage,
create_int64_atomic_all_ops_test(),
)
});
// See https://github.com/gfx-rs/wgpu/issues/5276 // See https://github.com/gfx-rs/wgpu/issues/5276
/* /*
fn create_int64_polyfill_test() -> Vec<ShaderTest> { fn create_int64_polyfill_test() -> Vec<ShaderTest> {

10
wgpu-core/src/device/mod.rs
View File

@ -491,6 +491,16 @@ pub fn create_validator(
Caps::SHADER_INT64, Caps::SHADER_INT64,
features.contains(wgt::Features::SHADER_INT64), features.contains(wgt::Features::SHADER_INT64),
); );
caps.set(
Caps::SHADER_INT64_ATOMIC_MIN_MAX,
features.intersects(
wgt::Features::SHADER_INT64_ATOMIC_MIN_MAX | wgt::Features::SHADER_INT64_ATOMIC_ALL_OPS,
),
);
caps.set(
Caps::SHADER_INT64_ATOMIC_ALL_OPS,
features.contains(wgt::Features::SHADER_INT64_ATOMIC_ALL_OPS),
);
caps.set( caps.set(
Caps::MULTISAMPLED_SHADING, Caps::MULTISAMPLED_SHADING,
downlevel.contains(wgt::DownlevelFlags::MULTISAMPLED_SHADING), downlevel.contains(wgt::DownlevelFlags::MULTISAMPLED_SHADING),

19
wgpu-hal/src/dx12/adapter.rs
View File

@ -354,6 +354,25 @@ impl super::Adapter {
&& features1.WaveOps != 0, && features1.WaveOps != 0,
); );
let atomic_int64_on_typed_resource_supported = {
let mut features9: crate::dx12::types::D3D12_FEATURE_DATA_D3D12_OPTIONS9 =
unsafe { mem::zeroed() };
let hr = unsafe {
device.CheckFeatureSupport(
37, // D3D12_FEATURE_D3D12_OPTIONS9
&mut features9 as *mut _ as *mut _,
mem::size_of::<crate::dx12::types::D3D12_FEATURE_DATA_D3D12_OPTIONS9>() as _,
)
};
hr == 0
&& features9.AtomicInt64OnGroupSharedSupported != 0
&& features9.AtomicInt64OnTypedResourceSupported != 0
};
features.set(
wgt::Features::SHADER_INT64_ATOMIC_ALL_OPS | wgt::Features::SHADER_INT64_ATOMIC_MIN_MAX,
atomic_int64_on_typed_resource_supported,
);
// float32-filterable should always be available on d3d12 // float32-filterable should always be available on d3d12
features.set(wgt::Features::FLOAT32_FILTERABLE, true); features.set(wgt::Features::FLOAT32_FILTERABLE, true);

18
wgpu-hal/src/dx12/types.rs
View File

@ -42,6 +42,24 @@ winapi::STRUCT! {
} }
} }
winapi::ENUM! {
enum D3D12_WAVE_MMA_TIER {
D3D12_WAVE_MMA_TIER_NOT_SUPPORTED = 0,
D3D12_WAVE_MMA_TIER_1_0 = 10,
}
}
winapi::STRUCT! {
struct D3D12_FEATURE_DATA_D3D12_OPTIONS9 {
MeshShaderPipelineStatsSupported: winapi::shared::minwindef::BOOL,
MeshShaderSupportsFullRangeRenderTargetArrayIndex: winapi::shared::minwindef::BOOL,
AtomicInt64OnTypedResourceSupported: winapi::shared::minwindef::BOOL,
AtomicInt64OnGroupSharedSupported: winapi::shared::minwindef::BOOL,
DerivativesInMeshAndAmplificationShadersSupported: winapi::shared::minwindef::BOOL,
WaveMMATier: D3D12_WAVE_MMA_TIER,
}
}
winapi::ENUM! { winapi::ENUM! {
enum D3D_SHADER_MODEL { enum D3D_SHADER_MODEL {
D3D_SHADER_MODEL_NONE = 0, D3D_SHADER_MODEL_NONE = 0,

9
wgpu-hal/src/metal/adapter.rs
View File

@ -821,6 +821,11 @@ impl super::PrivateCapabilities {
int64: family_check int64: family_check
&& (device.supports_family(MTLGPUFamily::Apple3) && (device.supports_family(MTLGPUFamily::Apple3)
|| device.supports_family(MTLGPUFamily::Metal3)), || device.supports_family(MTLGPUFamily::Metal3)),
// https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf#page=6
int64_atomics: family_check
&& ((device.supports_family(MTLGPUFamily::Apple8)
&& device.supports_family(MTLGPUFamily::Mac2))
|| device.supports_family(MTLGPUFamily::Apple9)),
} }
} }
@ -896,6 +901,10 @@ impl super::PrivateCapabilities {
F::SHADER_INT64, F::SHADER_INT64,
self.int64 && self.msl_version >= MTLLanguageVersion::V2_3, self.int64 && self.msl_version >= MTLLanguageVersion::V2_3,
); );
features.set(
F::SHADER_INT64_ATOMIC_MIN_MAX,
self.int64_atomics && self.msl_version >= MTLLanguageVersion::V2_4,
);
features.set( features.set(
F::ADDRESS_MODE_CLAMP_TO_BORDER, F::ADDRESS_MODE_CLAMP_TO_BORDER,

1
wgpu-hal/src/metal/mod.rs
View File

@ -272,6 +272,7 @@ struct PrivateCapabilities {
timestamp_query_support: TimestampQuerySupport, timestamp_query_support: TimestampQuerySupport,
supports_simd_scoped_operations: bool, supports_simd_scoped_operations: bool,
int64: bool, int64: bool,
int64_atomics: bool,
} }
#[derive(Clone, Debug)] #[derive(Clone, Debug)]

41
wgpu-hal/src/vulkan/adapter.rs
View File

@ -106,6 +106,9 @@ pub struct PhysicalDeviceFeatures {
zero_initialize_workgroup_memory: zero_initialize_workgroup_memory:
Option<vk::PhysicalDeviceZeroInitializeWorkgroupMemoryFeatures<'static>>, Option<vk::PhysicalDeviceZeroInitializeWorkgroupMemoryFeatures<'static>>,
/// Features provided by `VK_KHR_shader_atomic_int64`, promoted to Vulkan 1.2.
shader_atomic_int64: Option<vk::PhysicalDeviceShaderAtomicInt64Features<'static>>,
/// Features provided by `VK_EXT_subgroup_size_control`, promoted to Vulkan 1.3. /// Features provided by `VK_EXT_subgroup_size_control`, promoted to Vulkan 1.3.
subgroup_size_control: Option<vk::PhysicalDeviceSubgroupSizeControlFeatures<'static>>, subgroup_size_control: Option<vk::PhysicalDeviceSubgroupSizeControlFeatures<'static>>,
} }
@ -151,6 +154,9 @@ impl PhysicalDeviceFeatures {
if let Some(ref mut feature) = self.ray_query { if let Some(ref mut feature) = self.ray_query {
info = info.push_next(feature); info = info.push_next(feature);
} }
if let Some(ref mut feature) = self.shader_atomic_int64 {
info = info.push_next(feature);
}
if let Some(ref mut feature) = self.subgroup_size_control { if let Some(ref mut feature) = self.subgroup_size_control {
info = info.push_next(feature); info = info.push_next(feature);
} }
@ -419,6 +425,19 @@ impl PhysicalDeviceFeatures {
} else { } else {
None None
}, },
shader_atomic_int64: if device_api_version >= vk::API_VERSION_1_2
|| enabled_extensions.contains(&khr::shader_atomic_int64::NAME)
{
Some(
vk::PhysicalDeviceShaderAtomicInt64Features::default()
.shader_buffer_int64_atomics(requested_features.intersects(
wgt::Features::SHADER_INT64_ATOMIC_ALL_OPS
| wgt::Features::SHADER_INT64_ATOMIC_MIN_MAX,
)),
)
} else {
None
},
subgroup_size_control: if device_api_version >= vk::API_VERSION_1_3 subgroup_size_control: if device_api_version >= vk::API_VERSION_1_3
|| enabled_extensions.contains(&ext::subgroup_size_control::NAME) || enabled_extensions.contains(&ext::subgroup_size_control::NAME)
{ {
@ -559,6 +578,14 @@ impl PhysicalDeviceFeatures {
features.set(F::SHADER_INT64, self.core.shader_int64 != 0); features.set(F::SHADER_INT64, self.core.shader_int64 != 0);
features.set(F::SHADER_I16, self.core.shader_int16 != 0); features.set(F::SHADER_I16, self.core.shader_int16 != 0);
if let Some(ref shader_atomic_int64) = self.shader_atomic_int64 {
features.set(
F::SHADER_INT64_ATOMIC_ALL_OPS | F::SHADER_INT64_ATOMIC_MIN_MAX,
shader_atomic_int64.shader_buffer_int64_atomics != 0
&& shader_atomic_int64.shader_shared_int64_atomics != 0,
);
}
//if caps.supports_extension(khr::sampler_mirror_clamp_to_edge::NAME) { //if caps.supports_extension(khr::sampler_mirror_clamp_to_edge::NAME) {
//if caps.supports_extension(ext::sampler_filter_minmax::NAME) { //if caps.supports_extension(ext::sampler_filter_minmax::NAME) {
features.set( features.set(
@ -964,6 +991,13 @@ impl PhysicalDeviceProperties {
extensions.push(ext::texture_compression_astc_hdr::NAME); extensions.push(ext::texture_compression_astc_hdr::NAME);
} }
// Require `VK_KHR_shader_atomic_int64` if the associated feature was requested
if requested_features.intersects(
wgt::Features::SHADER_INT64_ATOMIC_ALL_OPS | wgt::Features::SHADER_INT64_ATOMIC_MIN_MAX,
) {
extensions.push(khr::shader_atomic_int64::NAME);
}
extensions extensions
} }
@ -1681,6 +1715,13 @@ impl super::Adapter {
capabilities.push(spv::Capability::Int64); capabilities.push(spv::Capability::Int64);
} }
if features.intersects(
wgt::Features::SHADER_INT64_ATOMIC_ALL_OPS
| wgt::Features::SHADER_INT64_ATOMIC_MIN_MAX,
) {
capabilities.push(spv::Capability::Int64Atomics);
}
let mut flags = spv::WriterFlags::empty(); let mut flags = spv::WriterFlags::empty();
flags.set( flags.set(
spv::WriterFlags::DEBUG, spv::WriterFlags::DEBUG,

17
wgpu-types/src/lib.rs
View File

@ -923,6 +923,23 @@ bitflags::bitflags! {
/// - DX12 /// - DX12
/// - Metal /// - Metal
const PIPELINE_CACHE = 1 << 59; const PIPELINE_CACHE = 1 << 59;
/// Allows shaders to use i64 and u64 atomic min and max.
///
/// Supported platforms:
/// - Vulkan (with VK_KHR_shader_atomic_int64)
/// - DX12 (with SM 6.6+)
/// - Metal (with MSL 2.4+)
///
/// This is a native only feature.
const SHADER_INT64_ATOMIC_MIN_MAX = 1 << 60;
/// Allows shaders to use all i64 and u64 atomic operations.
///
/// Supported platforms:
/// - Vulkan (with VK_KHR_shader_atomic_int64)
/// - DX12 (with SM 6.6+)
///
/// This is a native only feature.
const SHADER_INT64_ATOMIC_ALL_OPS = 1 << 61;
} }
} }