diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 679a2e3e1d..5fab23dce1 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -653,6 +653,7 @@ impl<'a, W: Write> Writer<'a, W> { match *inner { // Scalars are simple we just get the full name from `glsl_scalar` TypeInner::Scalar { kind, width } + | TypeInner::Atomic { kind, width } //TODO? | TypeInner::ValuePointer { size: None, kind, diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index ddaaf217a4..fbe59241fc 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -47,14 +47,15 @@ impl<'a> Display for TypeContext<'a> { match ty.inner { // work around Metal toolchain bug with `uint` typedef - crate::TypeInner::Scalar { - kind: crate::ScalarKind::Uint, - .. - } => { - write!(out, "metal::uint") - } crate::TypeInner::Scalar { kind, .. } => { - write!(out, "{}", scalar_kind_string(kind)) + let kind_str = match kind { + crate::ScalarKind::Uint => "metal::uint", + _ => scalar_kind_string(kind), + }; + write!(out, "{}", kind_str) + } + crate::TypeInner::Atomic { kind, .. } => { + write!(out, "atomic_{}", scalar_kind_string(kind)) } crate::TypeInner::Vector { size, kind, .. } => { write!( @@ -409,6 +410,7 @@ impl crate::Type { Ti::Scalar { .. } | Ti::Vector { .. } | Ti::Matrix { .. } + | Ti::Atomic { .. } | Ti::Pointer { .. } | Ti::ValuePointer { .. } => self.name.is_some(), // composite types are better to be aliased, regardless of the name diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs index e2ac8768c1..b390cbb5c0 100644 --- a/src/back/spv/writer.rs +++ b/src/back/spv/writer.rs @@ -671,7 +671,9 @@ impl Writer { use spirv::Decoration; let instruction = match ty.inner { - crate::TypeInner::Scalar { kind, width } => self.make_scalar(id, kind, width), + crate::TypeInner::Scalar { kind, width } | crate::TypeInner::Atomic { kind, width } => { + self.make_scalar(id, kind, width) + } crate::TypeInner::Vector { size, kind, width } => { let scalar_id = self.get_type_id(LookupType::Local(LocalType::Value { vector_size: None, diff --git a/src/front/wgsl/mod.rs b/src/front/wgsl/mod.rs index 54a4df5649..7fa77d0eea 100644 --- a/src/front/wgsl/mod.rs +++ b/src/front/wgsl/mod.rs @@ -347,7 +347,7 @@ impl<'a> Error<'a> { } impl crate::StorageFormat { - pub fn to_wgsl(self) -> &'static str { + fn to_wgsl(self) -> &'static str { use crate::StorageFormat as Sf; match self { Sf::R8Unorm => "r8unorm", @@ -392,17 +392,15 @@ impl crate::TypeInner { /// For example `vec3`. /// /// Note: The names of a `TypeInner::Struct` is not known. Therefore this method will simply return "struct" for them. - pub fn to_wgsl( - &self, - types: &Arena, - constants: &Arena, - ) -> String { + fn to_wgsl(&self, types: &Arena, constants: &Arena) -> String { + use crate::TypeInner as Ti; + match *self { - crate::TypeInner::Scalar { kind, width } => kind.to_wgsl(width), - crate::TypeInner::Vector { size, kind, width } => { + Ti::Scalar { kind, width } => kind.to_wgsl(width), + Ti::Vector { size, kind, width } => { format!("vec{}<{}>", size as u32, kind.to_wgsl(width)) } - crate::TypeInner::Matrix { + Ti::Matrix { columns, rows, width, @@ -414,15 +412,18 @@ impl crate::TypeInner { crate::ScalarKind::Float.to_wgsl(width), ) } - crate::TypeInner::Pointer { base, .. } => { + Ti::Atomic { kind, width } => { + format!("atomic<{}>", kind.to_wgsl(width)) + } + Ti::Pointer { base, .. } => { let base = &types[base]; let name = base.name.as_deref().unwrap_or("unknown"); - format!("*{}", name) + format!("ptr<{}>", name) } - crate::TypeInner::ValuePointer { kind, width, .. } => { - format!("*{}", kind.to_wgsl(width)) + Ti::ValuePointer { kind, width, .. } => { + format!("ptr<{}>", kind.to_wgsl(width)) } - crate::TypeInner::Array { base, size, .. } => { + Ti::Array { base, size, .. } => { let member_type = &types[base]; let base = member_type.name.as_deref().unwrap_or("unknown"); match size { @@ -433,11 +434,11 @@ impl crate::TypeInner { crate::ArraySize::Dynamic => format!("{}[]", base), } } - crate::TypeInner::Struct { .. } => { + Ti::Struct { .. } => { // TODO: Actually output the struct? "struct".to_string() } - crate::TypeInner::Image { + Ti::Image { dim, arrayed, class, @@ -481,7 +482,7 @@ impl crate::TypeInner { class_suffix, dim_suffix, array_suffix, type_in_brackets ) } - crate::TypeInner::Sampler { .. } => "sampler".to_string(), + Ti::Sampler { .. } => "sampler".to_string(), } } } @@ -537,7 +538,7 @@ mod type_inner_tests { access: crate::StorageAccess::default(), }, }; - assert_eq!(ptr.to_wgsl(&types, &constants), "*MyType2"); + assert_eq!(ptr.to_wgsl(&types, &constants), "ptr"); let img1 = crate::TypeInner::Image { dim: crate::ImageDimension::D2, diff --git a/src/lib.rs b/src/lib.rs index 469a5f496b..1e578662bf 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -509,6 +509,8 @@ pub enum TypeInner { rows: VectorSize, width: Bytes, }, + /// Atomic scalar. + Atomic { kind: ScalarKind, width: Bytes }, /// Pointer to another type. /// /// ## Pointers to non-`SIZED` types diff --git a/src/proc/layouter.rs b/src/proc/layouter.rs index 03ca286986..f2921f15e5 100644 --- a/src/proc/layouter.rs +++ b/src/proc/layouter.rs @@ -72,7 +72,7 @@ impl Layouter { for (ty_handle, ty) in types.iter().skip(self.layouts.len()) { let size = ty.inner.span(constants); let layout = match ty.inner { - Ti::Scalar { width, .. } => TypeLayout { + Ti::Scalar { width, .. } | Ti::Atomic { width, .. } => TypeLayout { size, alignment: Alignment::new(width as u32).unwrap(), }, diff --git a/src/proc/mod.rs b/src/proc/mod.rs index 2267808510..7919c4267e 100644 --- a/src/proc/mod.rs +++ b/src/proc/mod.rs @@ -95,7 +95,7 @@ impl super::TypeInner { pub fn span(&self, constants: &super::Arena) -> u32 { match *self { - Self::Scalar { kind: _, width } => width as u32, + Self::Scalar { kind: _, width } | Self::Atomic { kind: _, width } => width as u32, Self::Vector { size, kind: _, diff --git a/src/proc/typifier.rs b/src/proc/typifier.rs index 30af80d192..2978834980 100644 --- a/src/proc/typifier.rs +++ b/src/proc/typifier.rs @@ -350,7 +350,10 @@ impl<'a> ResolveContext<'a> { }) } crate::Expression::Load { pointer } => match *past(pointer).inner_with(types) { - Ti::Pointer { base, class: _ } => TypeResolution::Handle(base), + Ti::Pointer { base, class: _ } => match types[base].inner { + Ti::Atomic { kind, width } => TypeResolution::Value(Ti::Scalar { kind, width }), + _ => TypeResolution::Handle(base), + }, Ti::ValuePointer { size, kind, diff --git a/src/valid/mod.rs b/src/valid/mod.rs index d8b0be0d1f..f6d3d380b7 100644 --- a/src/valid/mod.rs +++ b/src/valid/mod.rs @@ -167,6 +167,7 @@ impl crate::TypeInner { size: crate::ArraySize::Constant(_), .. } + | Self::Atomic { .. } | Self::Pointer { .. } | Self::ValuePointer { .. } | Self::Struct { .. } => true, diff --git a/src/valid/type.rs b/src/valid/type.rs index f1ce40bce1..03869e80c1 100644 --- a/src/valid/type.rs +++ b/src/valid/type.rs @@ -35,20 +35,23 @@ bitflags::bitflags! { /// [`Struct`]: crate::Type::struct const SIZED = 0x2; + /// The data can be copied around. + const COPY = 0x4; + /// Can be be used for interfacing between pipeline stages. /// /// This includes non-bool scalars and vectors, matrices, and structs /// and arrays containing only interface types. - const INTERFACE = 0x4; + const INTERFACE = 0x8; /// Can be used for host-shareable structures. - const HOST_SHARED = 0x8; + const HOST_SHARED = 0x10; /// This is a top-level host-shareable type. - const TOP_LEVEL = 0x10; + const TOP_LEVEL = 0x20; /// This type can be passed as a function argument. - const ARGUMENT = 0x20; + const ARGUMENT = 0x40; } } @@ -72,6 +75,8 @@ pub enum Disalignment { pub enum TypeError { #[error("The {0:?} scalar width {1} is not supported")] InvalidWidth(crate::ScalarKind, crate::Bytes), + #[error("The {0:?} scalar width {1} is not supported for an atomic")] + InvalidAtomicWidth(crate::ScalarKind, crate::Bytes), #[error("The base handle {0:?} can not be resolved")] UnresolvedBase(Handle), #[error("Invalid type for pointer target {0:?}")] @@ -194,6 +199,7 @@ impl super::Validator { TypeInfo::new( TypeFlags::DATA | TypeFlags::SIZED + | TypeFlags::COPY | TypeFlags::INTERFACE | TypeFlags::HOST_SHARED | TypeFlags::ARGUMENT, @@ -208,6 +214,7 @@ impl super::Validator { TypeInfo::new( TypeFlags::DATA | TypeFlags::SIZED + | TypeFlags::COPY | TypeFlags::INTERFACE | TypeFlags::HOST_SHARED | TypeFlags::ARGUMENT, @@ -226,12 +233,26 @@ impl super::Validator { TypeInfo::new( TypeFlags::DATA | TypeFlags::SIZED + | TypeFlags::COPY | TypeFlags::INTERFACE | TypeFlags::HOST_SHARED | TypeFlags::ARGUMENT, count * (width as u32), ) } + Ti::Atomic { kind, width } => { + let good = match kind { + crate::ScalarKind::Bool | crate::ScalarKind::Float => false, + crate::ScalarKind::Sint | crate::ScalarKind::Uint => width == 4, + }; + if !good { + return Err(TypeError::InvalidAtomicWidth(kind, width)); + } + TypeInfo::new( + TypeFlags::DATA | TypeFlags::SIZED | TypeFlags::HOST_SHARED, + width as u32, + ) + } Ti::Pointer { base, class: _ } => { if base >= handle { return Err(TypeError::UnresolvedBase(base)); @@ -252,7 +273,7 @@ impl super::Validator { TypeFlags::empty() }; - TypeInfo::new(data_flag | TypeFlags::SIZED, 0) + TypeInfo::new(data_flag | TypeFlags::SIZED | TypeFlags::COPY, 0) } Ti::ValuePointer { size: _, @@ -263,7 +284,7 @@ impl super::Validator { if !self.check_width(kind, width) { return Err(TypeError::InvalidWidth(kind, width)); } - TypeInfo::new(TypeFlags::SIZED, 0) + TypeInfo::new(TypeFlags::DATA | TypeFlags::SIZED | TypeFlags::COPY, 0) } Ti::Array { base, size, stride } => { if base >= handle { @@ -366,7 +387,7 @@ impl super::Validator { } }; - let base_mask = TypeFlags::HOST_SHARED | TypeFlags::INTERFACE; + let base_mask = TypeFlags::COPY | TypeFlags::HOST_SHARED | TypeFlags::INTERFACE; TypeInfo { flags: TypeFlags::DATA | (base_info.flags & base_mask) | sized_flag, uniform_layout, @@ -381,6 +402,7 @@ impl super::Validator { let mut ti = TypeInfo::new( TypeFlags::DATA | TypeFlags::SIZED + | TypeFlags::COPY | TypeFlags::HOST_SHARED | TypeFlags::INTERFACE | TypeFlags::ARGUMENT,