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

Atomics in the IR

Browse Source
This commit is contained in:
Dzmitry Malyshau
2021-06-01 02:07:49 -04:00
committed by Dzmitry Malyshau
parent 551b711943
commit 8cb09c24c5
10 changed files with 70 additions and 36 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

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

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

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

4
src/back/spv/writer.rs
View File

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

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

@@ -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<f32>`.
///
/// 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<crate::Type>,
constants: &Arena<crate::Constant>,
) -> String {
fn to_wgsl(&self, types: &Arena<crate::Type>, constants: &Arena<crate::Constant>) -> 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<MyType2>");
let img1 = crate::TypeInner::Image {
dim: crate::ImageDimension::D2,

2
src/lib.rs
View File

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

2
src/proc/layouter.rs
View File

@@ -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(),
},

2
src/proc/mod.rs
View File

@@ -95,7 +95,7 @@ impl super::TypeInner {
pub fn span(&self, constants: &super::Arena<super::Constant>) -> 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: _,

5
src/proc/typifier.rs
View File

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

1
src/valid/mod.rs
View File

@@ -167,6 +167,7 @@ impl crate::TypeInner {
size: crate::ArraySize::Constant(_),
..
}
| Self::Atomic { .. }
| Self::Pointer { .. }
| Self::ValuePointer { .. }
| Self::Struct { .. } => true,

36
src/valid/type.rs
View File

@@ -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<crate::Type>),
#[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,