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

Enforce pointer semantics at the type level.

Browse Source
This commit is contained in:
Dzmitry Malyshau 2021-03-05 00:58:16 -05:00 committed by Dzmitry Malyshau
parent 7bc7288828
commit dd823a4b84
19 changed files with 934 additions and 566 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
Cargo.toml
View File

@ -22,6 +22,7 @@ thiserror = "1.0.21"
serde = { version = "1.0", features = ["derive"], optional = true }
petgraph = { version ="0.5", optional = true }
pp-rs = { git = "https://github.com/Kangz/glslpp-rs", rev = "4f2f72a", optional = true }
#env_logger = "0.8" # uncomment temporarily for developing with `convert`
[features]
default = []

2
bin/convert.rs
View File

@ -39,6 +39,8 @@ impl<T, E: Error> PrettyResult for Result<T, E> {
}
fn main() {
//env_logger::init(); // uncomment during development
let mut input_path = None;
let mut output_path = None;
//TODO: read the parameters from RON?

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

@ -538,11 +538,21 @@ impl<'a, W: Write> Writer<'a, W> {
fn write_type(&mut self, ty: Handle<Type>) -> BackendResult {
match self.module.types[ty].inner {
// Scalars are simple we just get the full name from `glsl_scalar`
TypeInner::Scalar { kind, width } => {
write!(self.out, "{}", glsl_scalar(kind, width)?.full)?
}
TypeInner::Scalar { kind, width }
| TypeInner::ValuePointer {
size: None,
kind,
width,
class: _,
} => write!(self.out, "{}", glsl_scalar(kind, width)?.full)?,
// Vectors are just `gvecN` where `g` is the scalar prefix and `N` is the vector size
TypeInner::Vector { size, kind, width } => write!(
TypeInner::Vector { size, kind, width }
| TypeInner::ValuePointer {
size: Some(size),
kind,
width,
class: _,
} => write!(
self.out,
"{}vec{}",
glsl_scalar(kind, width)?.prefix,
@ -1258,14 +1268,24 @@ impl<'a, W: Write> Writer<'a, W> {
Expression::AccessIndex { base, index } => {
self.write_expr(base, ctx)?;
match *ctx.typifier.get(base, &self.module.types) {
let mut resolved = ctx.typifier.get(base, &self.module.types);
let base_ty_handle = match *resolved {
TypeInner::Pointer { base, class: _ } => {
resolved = &self.module.types[base].inner;
Ok(base)
}
_ => ctx.typifier.get_handle(base),
};
match *resolved {
TypeInner::Vector { .. }
| TypeInner::Matrix { .. }
| TypeInner::Array { .. } => write!(self.out, "[{}]", index)?,
| TypeInner::Array { .. }
| TypeInner::ValuePointer { .. } => write!(self.out, "[{}]", index)?,
TypeInner::Struct { .. } => {
// This will never panic in case the type is a `Struct`, this is not true
// for other types so we can only check while inside this match arm
let ty = ctx.typifier.get_handle(base).unwrap();
let ty = base_ty_handle.unwrap();
write!(
self.out,

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

@ -118,6 +118,20 @@ impl crate::StorageClass {
_ => false,
}
}
fn get_name(&self, global_use: GlobalUse) -> Option<&'static str> {
match *self {
Self::Input | Self::Output | Self::Handle => None,
Self::Uniform => Some("constant"),
//TODO: should still be "constant" for read-only buffers
Self::Storage => Some(if global_use.contains(GlobalUse::WRITE) {
"device"
} else {
"storage "
}),
Self::Private | Self::Function | Self::WorkGroup | Self::PushConstant => Some(""),
}
}
}
enum FunctionOrigin {
@ -208,14 +222,21 @@ impl<W: Write> Writer<W> {
}
crate::Expression::AccessIndex { base, index } => {
self.put_expression(base, context)?;
let resolved = self.typifier.get(base, &context.module.types);
let mut resolved = self.typifier.get(base, &context.module.types);
let base_ty_handle = match *resolved {
crate::TypeInner::Pointer { base, class: _ } => {
resolved = &context.module.types[base].inner;
Ok(base)
}
_ => self.typifier.get_handle(base),
};
match *resolved {
crate::TypeInner::Struct { .. } => {
let base_ty = self.typifier.get_handle(base).unwrap();
let base_ty = base_ty_handle.unwrap();
let name = &self.names[&NameKey::StructMember(base_ty, index)];
write!(self.out, ".{}", name)?;
}
crate::TypeInner::Vector { .. } => {
crate::TypeInner::ValuePointer { .. } | crate::TypeInner::Vector { .. } => {
write!(self.out, ".{}", COMPONENTS[index as usize])?;
}
crate::TypeInner::Matrix { .. } => {
@ -852,6 +873,7 @@ impl<W: Write> Writer<W> {
fn write_type_defs(&mut self, module: &crate::Module) -> Result<(), Error> {
for (handle, ty) in module.types.iter() {
let name = &self.names[&NameKey::Type(handle)];
let global_use = GlobalUse::all(); //TODO
match ty.inner {
crate::TypeInner::Scalar { kind, .. } => {
write!(self.out, "typedef {} {}", scalar_kind_string(kind), name)?;
@ -878,20 +900,51 @@ impl<W: Write> Writer<W> {
)?;
}
crate::TypeInner::Pointer { base, class } => {
use crate::StorageClass as Sc;
let base_name = &self.names[&NameKey::Type(base)];
let class_name = match class {
Sc::Input | Sc::Output => continue,
Sc::Uniform => "constant",
Sc::Storage => "device",
Sc::Handle
| Sc::Private
| Sc::Function
| Sc::WorkGroup
| Sc::PushConstant => "",
let class_name = match class.get_name(global_use) {
Some(name) => name,
None => continue,
};
write!(self.out, "typedef {} {} *{}", class_name, base_name, name)?;
}
crate::TypeInner::ValuePointer {
size: None,
kind,
width: _,
class,
} => {
let class_name = match class.get_name(global_use) {
Some(name) => name,
None => continue,
};
write!(
self.out,
"typedef {} {} *{}",
class_name,
scalar_kind_string(kind),
name
)?;
}
crate::TypeInner::ValuePointer {
size: Some(size),
kind,
width: _,
class,
} => {
let class_name = match class.get_name(global_use) {
Some(name) => name,
None => continue,
};
write!(
self.out,
"typedef {} {}::{}{} {}",
class_name,
NAMESPACE,
scalar_kind_string(kind),
vector_size_string(size),
name
)?;
}
crate::TypeInner::Array {
base,
size,

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

@ -82,14 +82,11 @@ impl Function {
#[derive(Debug, PartialEq, Hash, Eq, Copy, Clone)]
enum LocalType {
Scalar {
kind: crate::ScalarKind,
width: crate::Bytes,
},
Vector {
size: crate::VectorSize,
Value {
vector_size: Option<crate::VectorSize>,
kind: crate::ScalarKind,
width: crate::Bytes,
pointer_class: Option<crate::StorageClass>,
},
Matrix {
columns: crate::VectorSize,
@ -108,10 +105,18 @@ enum LocalType {
impl LocalType {
fn from_inner(inner: &crate::TypeInner) -> Option<Self> {
Some(match *inner {
crate::TypeInner::Scalar { kind, width } => LocalType::Scalar { kind, width },
crate::TypeInner::Vector { size, kind, width } => {
LocalType::Vector { size, kind, width }
}
crate::TypeInner::Scalar { kind, width } => LocalType::Value {
vector_size: None,
kind,
width,
pointer_class: None,
},
crate::TypeInner::Vector { size, kind, width } => LocalType::Value {
vector_size: Some(size),
kind,
width,
pointer_class: None,
},
crate::TypeInner::Matrix {
columns,
rows,
@ -122,6 +127,17 @@ impl LocalType {
width,
},
crate::TypeInner::Pointer { base, class } => LocalType::Pointer { base, class },
crate::TypeInner::ValuePointer {
size,
kind,
width,
class,
} => LocalType::Value {
vector_size: size,
kind,
width,
pointer_class: Some(class),
},
_ => return None,
})
}
@ -302,26 +318,9 @@ impl Writer {
Ok(*e.get())
} else {
match lookup_ty {
LookupType::Handle(handle) => match arena[handle].inner {
crate::TypeInner::Scalar { kind, width } => self
.get_type_id(arena, LookupType::Local(LocalType::Scalar { kind, width })),
crate::TypeInner::Vector { size, kind, width } => self.get_type_id(
arena,
LookupType::Local(LocalType::Vector { size, kind, width }),
),
crate::TypeInner::Matrix {
columns,
rows,
width,
} => self.get_type_id(
arena,
LookupType::Local(LocalType::Matrix {
columns,
rows,
width,
}),
),
_ => self.write_type_declaration_arena(arena, handle),
LookupType::Handle(handle) => match LocalType::from_inner(&arena[handle].inner) {
Some(local) => self.get_type_id(arena, LookupType::Local(local)),
None => self.write_type_declaration_arena(arena, handle),
},
LookupType::Local(local_ty) => self.write_type_declaration_local(arena, local_ty),
}
@ -364,13 +363,6 @@ impl Writer {
)
}
fn create_pointer_type(&mut self, type_id: Word, class: spirv::StorageClass) -> Word {
let id = self.generate_id();
let instruction = Instruction::type_pointer(id, class, type_id);
instruction.to_words(&mut self.logical_layout.declarations);
id
}
fn create_constant(&mut self, type_id: Word, value: &[Word]) -> Word {
let id = self.generate_id();
let instruction = Instruction::constant(type_id, id, value);
@ -607,10 +599,27 @@ impl Writer {
) -> Result<Word, Error> {
let id = self.generate_id();
let instruction = match local_ty {
LocalType::Scalar { kind, width } => self.write_scalar(id, kind, width),
LocalType::Vector { size, kind, width } => {
let scalar_id =
self.get_type_id(arena, LookupType::Local(LocalType::Scalar { kind, width }))?;
LocalType::Value {
vector_size: None,
kind,
width,
pointer_class: None,
} => self.write_scalar(id, kind, width),
LocalType::Value {
vector_size: Some(size),
kind,
width,
pointer_class: None,
} => {
let scalar_id = self.get_type_id(
arena,
LookupType::Local(LocalType::Value {
vector_size: None,
kind,
width,
pointer_class: None,
}),
)?;
Instruction::type_vector(id, scalar_id, size)
}
LocalType::Matrix {
@ -620,16 +629,35 @@ impl Writer {
} => {
let vector_id = self.get_type_id(
arena,
LookupType::Local(LocalType::Vector {
size: rows,
LookupType::Local(LocalType::Value {
vector_size: Some(rows),
kind: crate::ScalarKind::Float,
width,
pointer_class: None,
}),
)?;
Instruction::type_matrix(id, vector_id, columns)
}
LocalType::Pointer { .. } => {
return Err(Error::FeatureNotImplemented("pointer declaration"))
LocalType::Pointer { base, class } => {
let type_id = self.get_type_id(arena, LookupType::Handle(base))?;
Instruction::type_pointer(id, self.parse_to_spirv_storage_class(class), type_id)
}
LocalType::Value {
vector_size,
kind,
width,
pointer_class: Some(class),
} => {
let type_id = self.get_type_id(
arena,
LookupType::Local(LocalType::Value {
vector_size,
kind,
width,
pointer_class: None,
}),
)?;
Instruction::type_pointer(id, self.parse_to_spirv_storage_class(class), type_id)
}
LocalType::SampledImage { image_type } => {
let image_type_id = self.get_type_id(arena, LookupType::Handle(image_type))?;
@ -658,15 +686,34 @@ impl Writer {
let instruction = match ty.inner {
crate::TypeInner::Scalar { kind, width } => {
self.lookup_type
.insert(LookupType::Local(LocalType::Scalar { kind, width }), id);
self.lookup_type.insert(
LookupType::Local(LocalType::Value {
vector_size: None,
kind,
width,
pointer_class: None,
}),
id,
);
self.write_scalar(id, kind, width)
}
crate::TypeInner::Vector { size, kind, width } => {
let scalar_id =
self.get_type_id(arena, LookupType::Local(LocalType::Scalar { kind, width }))?;
let scalar_id = self.get_type_id(
arena,
LookupType::Local(LocalType::Value {
vector_size: None,
kind,
width,
pointer_class: None,
}),
)?;
self.lookup_type.insert(
LookupType::Local(LocalType::Vector { size, kind, width }),
LookupType::Local(LocalType::Value {
vector_size: Some(size),
kind,
width,
pointer_class: None,
}),
id,
);
Instruction::type_vector(id, scalar_id, size)
@ -678,10 +725,11 @@ impl Writer {
} => {
let vector_id = self.get_type_id(
arena,
LookupType::Local(LocalType::Vector {
size: columns,
LookupType::Local(LocalType::Value {
vector_size: Some(columns),
kind: crate::ScalarKind::Float,
width,
pointer_class: None,
}),
)?;
self.lookup_type.insert(
@ -699,19 +747,16 @@ impl Writer {
arrayed,
class,
} => {
let width = 4;
let local_type = match class {
crate::ImageClass::Sampled { kind, multi: _ } => {
LocalType::Scalar { kind, width }
}
crate::ImageClass::Depth => LocalType::Scalar {
kind: crate::ScalarKind::Float,
width,
},
crate::ImageClass::Storage(format) => LocalType::Scalar {
kind: format.into(),
width,
},
let kind = match class {
crate::ImageClass::Sampled { kind, multi: _ } => kind,
crate::ImageClass::Depth => crate::ScalarKind::Float,
crate::ImageClass::Storage(format) => format.into(),
};
let local_type = LocalType::Value {
vector_size: None,
kind,
width: 4,
pointer_class: None,
};
let type_id = self.get_type_id(arena, LookupType::Local(local_type))?;
let dim = map_dim(dim);
@ -820,6 +865,32 @@ impl Writer {
.insert(LookupType::Local(LocalType::Pointer { base, class }), id);
Instruction::type_pointer(id, self.parse_to_spirv_storage_class(class), type_id)
}
crate::TypeInner::ValuePointer {
size,
kind,
width,
class,
} => {
let type_id = self.get_type_id(
arena,
LookupType::Local(LocalType::Value {
vector_size: size,
kind,
width,
pointer_class: None,
}),
)?;
self.lookup_type.insert(
LookupType::Local(LocalType::Value {
vector_size: size,
kind,
width,
pointer_class: Some(class),
}),
id,
);
Instruction::type_pointer(id, self.parse_to_spirv_storage_class(class), type_id)
}
};
self.lookup_type.insert(LookupType::Handle(handle), id);
@ -837,9 +908,11 @@ impl Writer {
crate::ConstantInner::Scalar { width, ref value } => {
let type_id = self.get_type_id(
types,
LookupType::Local(LocalType::Scalar {
LookupType::Local(LocalType::Value {
vector_size: None,
kind: value.scalar_kind(),
width,
pointer_class: None,
}),
)?;
let (solo, pair);
@ -1061,9 +1134,11 @@ impl Writer {
Ok(if let Some(array_index) = array_index {
let coordinate_scalar_type_id = self.get_type_id(
&ir_module.types,
LookupType::Local(LocalType::Scalar {
LookupType::Local(LocalType::Value {
vector_size: None,
kind: crate::ScalarKind::Float,
width: 4,
pointer_class: None,
}),
)?;
@ -1109,10 +1184,11 @@ impl Writer {
let extended_coordinate_type_id = self.get_type_id(
&ir_module.types,
LookupType::Local(LocalType::Vector {
size,
LookupType::Local(LocalType::Value {
vector_size: Some(size),
kind: crate::ScalarKind::Float,
width: 4,
pointer_class: None,
}),
)?;
@ -1757,6 +1833,7 @@ impl Writer {
Err(inner) => LookupType::Local(LocalType::from_inner(inner).unwrap()),
};
let result_type_id = self.get_type_id(&ir_module.types, result_lookup_ty)?;
self.temp_chain.clear();
let (root_id, class) = loop {
expr_handle = match ir_function.expressions[expr_handle] {
@ -1768,9 +1845,11 @@ impl Writer {
crate::Expression::AccessIndex { base, index } => {
let const_ty_id = self.get_type_id(
&ir_module.types,
LookupType::Local(LocalType::Scalar {
LookupType::Local(LocalType::Value {
vector_size: None,
kind: crate::ScalarKind::Sint,
width: 4,
pointer_class: None,
}),
)?;
let const_id = self.create_constant(const_ty_id, &[index]);
@ -1794,9 +1873,8 @@ impl Writer {
} else {
self.temp_chain.reverse();
let id = self.generate_id();
let pointer_type_id = self.create_pointer_type(result_type_id, class);
block.body.push(Instruction::access_chain(
pointer_type_id,
result_type_id,
id,
root_id,
&self.temp_chain,

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

@ -136,6 +136,8 @@ pub enum Error<'a> {
InvalidResolve(ResolveError),
#[error("invalid statement {0:?}, expected {1}")]
InvalidStatement(crate::Statement, &'a str),
#[error("resource type {0:?} is invalid")]
InvalidResourceType(Handle<crate::Type>),
#[error("unknown import: `{0}`")]
UnknownImport(&'a str),
#[error("unknown storage class: `{0}`")]
@ -170,6 +172,8 @@ pub enum Error<'a> {
FunctionRedefinition(&'a str),
#[error("call to local `{0}(..)` can't be resolved")]
UnknownLocalFunction(&'a str),
#[error("builtin {0:?} is not implemented")]
UnimplementedBuiltin(crate::BuiltIn),
#[error("other error")]
Other,
}
@ -343,6 +347,23 @@ impl Composition {
}
}
fn extract(
base: Handle<crate::Expression>,
base_size: crate::VectorSize,
name: &str,
name_span: Range<usize>,
) -> Result<crate::Expression, Error> {
let ch = name
.chars()
.next()
.ok_or_else(|| Error::BadAccessor(name_span.clone()))?;
let index = Self::letter_pos(ch);
if index >= base_size as u32 {
return Err(Error::BadAccessor(name_span));
}
Ok(crate::Expression::AccessIndex { base, index })
}
fn make<'a>(
base: Handle<crate::Expression>,
base_size: crate::VectorSize,
@ -350,7 +371,7 @@ impl Composition {
name_span: Range<usize>,
expressions: &mut Arena<crate::Expression>,
) -> Result<Self, Error<'a>> {
Ok(if name.len() > 1 {
if name.len() > 1 {
let mut components = Vec::with_capacity(name.len());
for ch in name.chars() {
let index = Self::letter_pos(ch);
@ -367,18 +388,10 @@ impl Composition {
4 => crate::VectorSize::Quad,
_ => return Err(Error::BadAccessor(name_span)),
};
Composition::Multi(size, components)
Ok(Composition::Multi(size, components))
} else {
let ch = name
.chars()
.next()
.ok_or_else(|| Error::BadAccessor(name_span.clone()))?;
let index = Self::letter_pos(ch);
if index >= base_size as u32 {
return Err(Error::BadAccessor(name_span));
}
Composition::Single(crate::Expression::AccessIndex { base, index })
})
Self::extract(base, base_size, name, name_span).map(Composition::Single)
}
}
}
@ -1070,10 +1083,19 @@ impl Parser {
};
loop {
// insert the E::Load when we reach a value
if needs_deref && ctx.resolve_type(handle)?.scalar_kind().is_some() {
let expression = crate::Expression::Load { pointer: handle };
handle = ctx.expressions.append(expression);
needs_deref = false;
if needs_deref {
let now = match *ctx.resolve_type(handle)? {
crate::TypeInner::Pointer { base, class: _ } => {
ctx.types[base].inner.scalar_kind().is_some()
}
crate::TypeInner::ValuePointer { .. } => true,
_ => false,
};
if now {
let expression = crate::Expression::Load { pointer: handle };
handle = ctx.expressions.append(expression);
needs_deref = false;
}
}
let expression = match lexer.peek().0 {
@ -1133,6 +1155,30 @@ impl Parser {
}
Composition::Single(expr) => expr,
},
crate::TypeInner::ValuePointer {
size: Some(size), ..
} => Composition::extract(handle, size, name, name_span)?,
crate::TypeInner::Pointer { base, class: _ } => match ctx.types[base].inner
{
crate::TypeInner::Struct { ref members, .. } => {
let index = members
.iter()
.position(|m| m.name.as_deref() == Some(name))
.ok_or(Error::BadAccessor(name_span))?
as u32;
crate::Expression::AccessIndex {
base: handle,
index,
}
}
crate::TypeInner::Vector { size, .. } => {
Composition::extract(handle, size, name, name_span)?
}
crate::TypeInner::Matrix { columns, .. } => {
Composition::extract(handle, columns, name, name_span)?
}
_ => return Err(Error::BadAccessor(name_span)),
},
_ => return Err(Error::BadAccessor(name_span)),
}
}
@ -2453,18 +2499,29 @@ impl Parser {
Some(crate::Binding::BuiltIn(builtin)) => match builtin {
crate::BuiltIn::GlobalInvocationId => crate::StorageClass::Input,
crate::BuiltIn::Position => crate::StorageClass::Output,
_ => unimplemented!(),
_ => return Err(Error::UnimplementedBuiltin(builtin)),
},
Some(crate::Binding::Resource { .. }) => {
match module.types[pvar.ty].inner {
crate::TypeInner::Struct { .. } if pvar.access.is_empty() => {
crate::StorageClass::Uniform
}
crate::TypeInner::Struct { .. } => crate::StorageClass::Storage,
_ => crate::StorageClass::Handle,
crate::TypeInner::Struct { .. }
| crate::TypeInner::Array { .. } => crate::StorageClass::Storage,
crate::TypeInner::Image { .. }
| crate::TypeInner::Sampler { .. } => crate::StorageClass::Handle,
ref other => {
log::error!("Resource type {:?}", other);
return Err(Error::InvalidResourceType(pvar.ty));
}
}
}
_ => crate::StorageClass::Private,
_ => match module.types[pvar.ty].inner {
crate::TypeInner::Image { .. } | crate::TypeInner::Sampler { .. } => {
crate::StorageClass::Handle
}
_ => crate::StorageClass::Private,
},
},
};
let var_handle = module.global_variables.append(crate::GlobalVariable {

12
src/lib.rs
View File

@ -53,6 +53,9 @@ use serde::Deserialize;
#[cfg(feature = "serialize")]
use serde::Serialize;
/// Width of a boolean type, in bytes.
pub const BOOL_WIDTH: Bytes = 1;
/// Hash map that is faster but not resilient to DoS attacks.
pub type FastHashMap<K, T> = HashMap<K, T, BuildHasherDefault<fxhash::FxHasher>>;
/// Hash set that is faster but not resilient to DoS attacks.
@ -371,11 +374,18 @@ pub enum TypeInner {
rows: VectorSize,
width: Bytes,
},
/// Pointer to a value.
/// Pointer to another type.
Pointer {
base: Handle<Type>,
class: StorageClass,
},
/// Pointer to a value.
ValuePointer {
size: Option<VectorSize>,
kind: ScalarKind,
width: Bytes,
class: StorageClass,
},
/// Homogenous list of elements.
Array {
base: Handle<Type>,

2
src/proc/layouter.rs
View File

@ -62,7 +62,7 @@ impl Layouter {
size: (columns as u8 * rows as u8 * width) as u32,
alignment: Alignment::new((columns as u8 * width) as u32).unwrap(),
},
Ti::Pointer { .. } => TypeLayout {
Ti::Pointer { .. } | Ti::ValuePointer { .. } => TypeLayout {
size: 4,
alignment: Alignment::new(1).unwrap(),
},

332
src/proc/typifier.rs
View File

@ -11,26 +11,34 @@ enum Resolution {
// Clone is only implemented for numeric variants of `TypeInner`.
impl Clone for Resolution {
fn clone(&self) -> Self {
use crate::TypeInner as Ti;
match *self {
Resolution::Handle(handle) => Resolution::Handle(handle),
Resolution::Value(ref v) => Resolution::Value(match *v {
crate::TypeInner::Scalar { kind, width } => {
crate::TypeInner::Scalar { kind, width }
}
crate::TypeInner::Vector { size, kind, width } => {
crate::TypeInner::Vector { size, kind, width }
}
crate::TypeInner::Matrix {
Ti::Scalar { kind, width } => Ti::Scalar { kind, width },
Ti::Vector { size, kind, width } => Ti::Vector { size, kind, width },
Ti::Matrix {
rows,
columns,
width,
} => crate::TypeInner::Matrix {
} => Ti::Matrix {
rows,
columns,
width,
},
#[allow(clippy::panic)]
_ => panic!("Unexpected clone type: {:?}", v),
Ti::Pointer { base, class } => Ti::Pointer { base, class },
Ti::ValuePointer {
size,
kind,
width,
class,
} => Ti::ValuePointer {
size,
kind,
width,
class,
},
_ => unreachable!("Unexpected clone type: {:?}", v),
}),
}
}
@ -44,8 +52,25 @@ pub struct Typifier {
#[derive(Clone, Debug, Error, PartialEq)]
pub enum ResolveError {
#[error("Invalid index into array")]
InvalidAccessIndex,
#[error("Index {index} is out of bounds for expression {expr:?}")]
OutOfBoundsIndex {
expr: Handle<crate::Expression>,
index: u32,
},
#[error("Invalid access into expression {expr:?}, indexed: {indexed}")]
InvalidAccess {
expr: Handle<crate::Expression>,
indexed: bool,
},
#[error("Invalid sub-access into type {ty:?}, indexed: {indexed}")]
InvalidSubAccess {
ty: Handle<crate::Type>,
indexed: bool,
},
#[error("Invalid pointer {0:?}")]
InvalidPointer(Handle<crate::Expression>),
#[error("Invalid image {0:?}")]
InvalidImage(Handle<crate::Expression>),
#[error("Function {name} not defined")]
FunctionNotDefined { name: String },
#[error("Function without return type")]
@ -118,51 +143,74 @@ impl Typifier {
}
}
//TODO: resolve `*Variable` and `Access*` expressions to `Pointer` type.
fn resolve_impl(
&self,
expr: &crate::Expression,
types: &Arena<crate::Type>,
ctx: &ResolveContext,
) -> Result<Resolution, ResolveError> {
use crate::TypeInner as Ti;
Ok(match *expr {
crate::Expression::Access { base, .. } => match *self.get(base, types) {
crate::TypeInner::Array { base, .. } => Resolution::Handle(base),
crate::TypeInner::Vector {
Ti::Array { base, .. } => Resolution::Handle(base),
Ti::Vector {
size: _,
kind,
width,
} => Resolution::Value(crate::TypeInner::Scalar { kind, width }),
crate::TypeInner::Matrix {
rows: size,
columns: _,
} => Resolution::Value(Ti::Scalar { kind, width }),
Ti::ValuePointer {
size: Some(_),
kind,
width,
} => Resolution::Value(crate::TypeInner::Vector {
size,
kind: crate::ScalarKind::Float,
class,
} => Resolution::Value(Ti::ValuePointer {
size: None,
kind,
width,
class,
}),
Ti::Pointer { base, class } => Resolution::Value(match types[base].inner {
Ti::Array { base, .. } => Ti::Pointer { base, class },
Ti::Vector {
size: _,
kind,
width,
} => Ti::ValuePointer {
size: None,
kind,
width,
class,
},
ref other => {
log::error!("Access sub-type {:?}", other);
return Err(ResolveError::InvalidSubAccess {
ty: base,
indexed: false,
});
}
}),
ref other => {
return Err(ResolveError::IncompatibleOperand {
op: "access".to_string(),
operand: format!("{:?}", other),
})
log::error!("Access type {:?}", other);
return Err(ResolveError::InvalidAccess {
expr: base,
indexed: false,
});
}
},
crate::Expression::AccessIndex { base, index } => match *self.get(base, types) {
crate::TypeInner::Vector { size, kind, width } => {
Ti::Vector { size, kind, width } => {
if index >= size as u32 {
return Err(ResolveError::InvalidAccessIndex);
return Err(ResolveError::OutOfBoundsIndex { expr: base, index });
}
Resolution::Value(crate::TypeInner::Scalar { kind, width })
Resolution::Value(Ti::Scalar { kind, width })
}
crate::TypeInner::Matrix {
Ti::Matrix {
columns,
rows,
width,
} => {
if index >= columns as u32 {
return Err(ResolveError::InvalidAccessIndex);
return Err(ResolveError::OutOfBoundsIndex { expr: base, index });
}
Resolution::Value(crate::TypeInner::Vector {
size: rows,
@ -170,26 +218,94 @@ impl Typifier {
width,
})
}
crate::TypeInner::Array { base, .. } => Resolution::Handle(base),
crate::TypeInner::Struct {
Ti::Array { base, .. } => Resolution::Handle(base),
Ti::Struct {
block: _,
ref members,
} => {
let member = members
.get(index as usize)
.ok_or(ResolveError::InvalidAccessIndex)?;
.ok_or(ResolveError::OutOfBoundsIndex { expr: base, index })?;
Resolution::Handle(member.ty)
}
ref other => {
return Err(ResolveError::IncompatibleOperand {
op: "access index".to_string(),
operand: format!("{:?}", other),
Ti::ValuePointer {
size: Some(size),
kind,
width,
class,
} => {
if index >= size as u32 {
return Err(ResolveError::OutOfBoundsIndex { expr: base, index });
}
Resolution::Value(Ti::ValuePointer {
size: None,
kind,
width,
class,
})
}
Ti::Pointer {
base: ty_base,
class,
} => Resolution::Value(match types[ty_base].inner {
Ti::Array { base, .. } => Ti::Pointer { base, class },
Ti::Vector { size, kind, width } => {
if index >= size as u32 {
return Err(ResolveError::OutOfBoundsIndex { expr: base, index });
}
Ti::ValuePointer {
size: None,
kind,
width,
class,
}
}
Ti::Matrix {
rows,
columns,
width,
} => {
if index >= columns as u32 {
return Err(ResolveError::OutOfBoundsIndex { expr: base, index });
}
Ti::ValuePointer {
size: Some(rows),
kind: crate::ScalarKind::Float,
width,
class,
}
}
Ti::Struct {
block: _,
ref members,
} => {
let member = members
.get(index as usize)
.ok_or(ResolveError::OutOfBoundsIndex { expr: base, index })?;
Ti::Pointer {
base: member.ty,
class,
}
}
ref other => {
log::error!("Access index sub-type {:?}", other);
return Err(ResolveError::InvalidSubAccess {
ty: ty_base,
indexed: true,
});
}
}),
ref other => {
log::error!("Access index type {:?}", other);
return Err(ResolveError::InvalidAccess {
expr: base,
indexed: true,
});
}
},
crate::Expression::Constant(h) => match ctx.constants[h].inner {
crate::ConstantInner::Scalar { width, ref value } => {
Resolution::Value(crate::TypeInner::Scalar {
Resolution::Value(Ti::Scalar {
kind: value.scalar_kind(),
width,
})
@ -200,55 +316,89 @@ impl Typifier {
crate::Expression::FunctionArgument(index) => {
Resolution::Handle(ctx.arguments[index as usize].ty)
}
crate::Expression::GlobalVariable(h) => Resolution::Handle(ctx.global_vars[h].ty),
crate::Expression::LocalVariable(h) => Resolution::Handle(ctx.local_vars[h].ty),
// we treat Load as a transparent operation for the type system
crate::Expression::Load { pointer } => self.resolutions[pointer.index()].clone(),
crate::Expression::GlobalVariable(h) => {
let var = &ctx.global_vars[h];
if var.class == crate::StorageClass::Handle {
Resolution::Handle(var.ty)
} else {
Resolution::Value(Ti::Pointer {
base: var.ty,
class: var.class,
})
}
}
crate::Expression::LocalVariable(h) => {
let var = &ctx.local_vars[h];
Resolution::Value(Ti::Pointer {
base: var.ty,
class: crate::StorageClass::Function,
})
}
crate::Expression::Load { pointer } => match *self.get(pointer, types) {
Ti::Pointer { base, class: _ } => Resolution::Handle(base),
Ti::ValuePointer {
size,
kind,
width,
class: _,
} => Resolution::Value(match size {
Some(size) => Ti::Vector { size, kind, width },
None => Ti::Scalar { kind, width },
}),
ref other => {
log::error!("Pointer type {:?}", other);
return Err(ResolveError::InvalidPointer(pointer));
}
},
crate::Expression::ImageSample { image, .. }
| crate::Expression::ImageLoad { image, .. } => match *self.get(image, types) {
crate::TypeInner::Image { class, .. } => Resolution::Value(match class {
crate::ImageClass::Depth => crate::TypeInner::Scalar {
Ti::Image { class, .. } => Resolution::Value(match class {
crate::ImageClass::Depth => Ti::Scalar {
kind: crate::ScalarKind::Float,
width: 4,
},
crate::ImageClass::Sampled { kind, multi: _ } => crate::TypeInner::Vector {
crate::ImageClass::Sampled { kind, multi: _ } => Ti::Vector {
kind,
width: 4,
size: crate::VectorSize::Quad,
},
crate::ImageClass::Storage(format) => crate::TypeInner::Vector {
crate::ImageClass::Storage(format) => Ti::Vector {
kind: format.into(),
width: 4,
size: crate::VectorSize::Quad,
},
}),
_ => unreachable!(),
ref other => {
log::error!("Image type {:?}", other);
return Err(ResolveError::InvalidImage(image));
}
},
crate::Expression::ImageQuery { image, query } => Resolution::Value(match query {
crate::ImageQuery::Size { level: _ } => match *self.get(image, types) {
crate::TypeInner::Image { dim, .. } => match dim {
crate::ImageDimension::D1 => crate::TypeInner::Scalar {
Ti::Image { dim, .. } => match dim {
crate::ImageDimension::D1 => Ti::Scalar {
kind: crate::ScalarKind::Sint,
width: 4,
},
crate::ImageDimension::D2 => crate::TypeInner::Vector {
crate::ImageDimension::D2 => Ti::Vector {
size: crate::VectorSize::Bi,
kind: crate::ScalarKind::Sint,
width: 4,
},
crate::ImageDimension::D3 | crate::ImageDimension::Cube => {
crate::TypeInner::Vector {
size: crate::VectorSize::Tri,
kind: crate::ScalarKind::Sint,
width: 4,
}
}
crate::ImageDimension::D3 | crate::ImageDimension::Cube => Ti::Vector {
size: crate::VectorSize::Tri,
kind: crate::ScalarKind::Sint,
width: 4,
},
},
_ => unreachable!(),
ref other => {
log::error!("Image type {:?}", other);
return Err(ResolveError::InvalidImage(image));
}
},
crate::ImageQuery::NumLevels
| crate::ImageQuery::NumLayers
| crate::ImageQuery::NumSamples => crate::TypeInner::Scalar {
| crate::ImageQuery::NumSamples => Ti::Scalar {
kind: crate::ScalarKind::Sint,
width: 4,
},
@ -264,28 +414,28 @@ impl Typifier {
let ty_right = self.get(right, types);
if ty_left == ty_right {
self.resolutions[left.index()].clone()
} else if let crate::TypeInner::Scalar { .. } = *ty_left {
} else if let Ti::Scalar { .. } = *ty_left {
self.resolutions[right.index()].clone()
} else if let crate::TypeInner::Scalar { .. } = *ty_right {
} else if let Ti::Scalar { .. } = *ty_right {
self.resolutions[left.index()].clone()
} else if let crate::TypeInner::Matrix {
} else if let Ti::Matrix {
columns: _,
rows,
width,
} = *ty_left
{
Resolution::Value(crate::TypeInner::Vector {
Resolution::Value(Ti::Vector {
size: rows,
kind: crate::ScalarKind::Float,
width,
})
} else if let crate::TypeInner::Matrix {
} else if let Ti::Matrix {
columns,
rows: _,
width,
} = *ty_right
{
Resolution::Value(crate::TypeInner::Vector {
Resolution::Value(Ti::Vector {
size: columns,
kind: crate::ScalarKind::Float,
width,
@ -309,10 +459,8 @@ impl Typifier {
let kind = crate::ScalarKind::Bool;
let width = 1;
let inner = match *self.get(left, types) {
crate::TypeInner::Scalar { .. } => crate::TypeInner::Scalar { kind, width },
crate::TypeInner::Vector { size, .. } => {
crate::TypeInner::Vector { size, kind, width }
}
Ti::Scalar { .. } => Ti::Scalar { kind, width },
Ti::Vector { size, .. } => Ti::Vector { size, kind, width },
ref other => {
return Err(ResolveError::IncompatibleOperand {
op: "logical".to_string(),
@ -332,7 +480,7 @@ impl Typifier {
crate::Expression::Derivative { axis: _, expr } => {
self.resolutions[expr.index()].clone()
}
crate::Expression::Relational { .. } => Resolution::Value(crate::TypeInner::Scalar {
crate::Expression::Relational { .. } => Resolution::Value(Ti::Scalar {
kind: crate::ScalarKind::Bool,
width: 4,
}),
@ -377,11 +525,11 @@ impl Typifier {
Mf::Pow => self.resolutions[arg.index()].clone(),
// geometry
Mf::Dot => match *self.get(arg, types) {
crate::TypeInner::Vector {
Ti::Vector {
kind,
size: _,
width,
} => Resolution::Value(crate::TypeInner::Scalar { kind, width }),
} => Resolution::Value(Ti::Scalar { kind, width }),
ref other => {
return Err(ResolveError::IncompatibleOperand {
op: "dot product".to_string(),
@ -395,7 +543,7 @@ impl Typifier {
operand: "".to_string(),
})?;
match (self.get(arg, types), self.get(arg1,types)) {
(&crate::TypeInner::Vector {kind: _, size: columns,width}, &crate::TypeInner::Vector{ size: rows, .. }) => Resolution::Value(crate::TypeInner::Matrix { columns, rows, width }),
(&Ti::Vector {kind: _, size: columns,width}, &Ti::Vector{ size: rows, .. }) => Resolution::Value(Ti::Matrix { columns, rows, width }),
(left, right) => {
return Err(ResolveError::IncompatibleOperands {
op: "outer product".to_string(),
@ -408,8 +556,8 @@ impl Typifier {
Mf::Cross => self.resolutions[arg.index()].clone(),
Mf::Distance |
Mf::Length => match *self.get(arg, types) {
crate::TypeInner::Scalar {width,kind} |
crate::TypeInner::Vector {width,kind,size:_} => Resolution::Value(crate::TypeInner::Scalar { kind, width }),
Ti::Scalar {width,kind} |
Ti::Vector {width,kind,size:_} => Resolution::Value(Ti::Scalar { kind, width }),
ref other => {
return Err(ResolveError::IncompatibleOperand {
op: format!("{:?}", fun),
@ -429,11 +577,11 @@ impl Typifier {
Mf::Sqrt |
Mf::InverseSqrt => self.resolutions[arg.index()].clone(),
Mf::Transpose => match *self.get(arg, types) {
crate::TypeInner::Matrix {
Ti::Matrix {
columns,
rows,
width,
} => Resolution::Value(crate::TypeInner::Matrix {
} => Resolution::Value(Ti::Matrix {
columns: rows,
rows: columns,
width,
@ -446,11 +594,11 @@ impl Typifier {
}
},
Mf::Inverse => match *self.get(arg, types) {
crate::TypeInner::Matrix {
Ti::Matrix {
columns,
rows,
width,
} if columns == rows => Resolution::Value(crate::TypeInner::Matrix {
} if columns == rows => Resolution::Value(Ti::Matrix {
columns,
rows,
width,
@ -463,10 +611,10 @@ impl Typifier {
}
},
Mf::Determinant => match *self.get(arg, types) {
crate::TypeInner::Matrix {
Ti::Matrix {
width,
..
} => Resolution::Value(crate::TypeInner::Scalar { kind: crate::ScalarKind::Float, width }),
} => Resolution::Value(Ti::Scalar { kind: crate::ScalarKind::Float, width }),
ref other => {
return Err(ResolveError::IncompatibleOperand {
op: "determinant".to_string(),
@ -484,14 +632,12 @@ impl Typifier {
kind,
convert: _,
} => match *self.get(expr, types) {
crate::TypeInner::Scalar { kind: _, width } => {
Resolution::Value(crate::TypeInner::Scalar { kind, width })
}
crate::TypeInner::Vector {
Ti::Scalar { kind: _, width } => Resolution::Value(Ti::Scalar { kind, width }),
Ti::Vector {
kind: _,
size,
width,
} => Resolution::Value(crate::TypeInner::Vector { kind, size, width }),
} => Resolution::Value(Ti::Vector { kind, size, width }),
ref other => {
return Err(ResolveError::IncompatibleOperand {
op: "as".to_string(),
@ -505,7 +651,7 @@ impl Typifier {
.ok_or(ResolveError::FunctionReturnsVoid)?;
Resolution::Handle(ty)
}
crate::Expression::ArrayLength(_) => Resolution::Value(crate::TypeInner::Scalar {
crate::Expression::ArrayLength(_) => Resolution::Value(Ti::Scalar {
kind: crate::ScalarKind::Uint,
width: 4,
}),
@ -523,15 +669,7 @@ impl Typifier {
for (eh, expr) in expressions.iter().skip(self.resolutions.len()) {
let resolution = self.resolve_impl(expr, types, ctx)?;
log::debug!("Resolving {:?} = {:?} : {:?}", eh, expr, resolution);
let ty_handle = match resolution {
Resolution::Handle(h) => h,
Resolution::Value(inner) => types
.fetch_if_or_append(crate::Type { name: None, inner }, |a, b| {
a.inner == b.inner
}),
};
self.resolutions.push(Resolution::Handle(ty_handle));
self.resolutions.push(resolution);
}
}
Ok(())

33
src/proc/validator.rs
View File

@ -455,7 +455,7 @@ impl Validator {
fn check_width(kind: crate::ScalarKind, width: crate::Bytes) -> bool {
match kind {
crate::ScalarKind::Bool => width == 1,
crate::ScalarKind::Bool => width == crate::BOOL_WIDTH,
_ => width == 4,
}
}
@ -486,6 +486,17 @@ impl Validator {
}
TypeFlags::DATA | TypeFlags::SIZED
}
Ti::ValuePointer {
size: _,
kind,
width,
class: _,
} => {
if !Self::check_width(kind, width) {
return Err(TypeError::InvalidWidth(kind, width));
}
TypeFlags::SIZED //TODO: `DATA`?
}
Ti::Array { base, size, stride } => {
if base >= handle {
return Err(TypeError::UnresolvedBase(base));
@ -917,7 +928,25 @@ impl Validator {
}
_ => {}
}
if self.typifier.try_get(pointer, context.types) != Some(value_ty) {
let good = match self.typifier.try_get(pointer, context.types) {
Some(&Ti::Pointer { base, class: _ }) => {
*value_ty == context.types[base].inner
}
Some(&Ti::ValuePointer {
size: Some(size),
kind,
width,
class: _,
}) => *value_ty == Ti::Vector { size, kind, width },
Some(&Ti::ValuePointer {
size: None,
kind,
width,
class: _,
}) => *value_ty == Ti::Scalar { kind, width },
_ => false,
};
if !good {
return Err(FunctionError::InvalidStoreTypes { pointer, value });
}
}

2
tests/out/boids.msl.snap
View File

@ -36,8 +36,6 @@ typedef metal::uint3 type4;
typedef int type5;
typedef bool type6;
constexpr constant int NUM_PARTICLES = 1500;
constexpr constant float const_0f = 0.0;
constexpr constant int const_0i = 0;

440
tests/out/boids.spvasm.snap
View File

@ -5,7 +5,7 @@ expression: dis
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 233
; Bound: 221
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
@ -94,46 +94,34 @@ OpDecorate %25 BuiltIn GlobalInvocationId
%40 = OpTypePointer Function %9
%42 = OpTypeFunction %2
%47 = OpTypeBool
%51 = OpConstant %4 0
%52 = OpConstant %4 0
%54 = OpTypePointer Uniform %22
%56 = OpConstant %4 1
%57 = OpConstant %4 0
%59 = OpTypePointer Uniform %22
%77 = OpConstant %4 0
%51 = OpTypePointer Uniform %20
%52 = OpTypePointer Uniform %21
%53 = OpTypePointer Uniform %22
%54 = OpConstant %4 0
%55 = OpConstant %4 0
%58 = OpConstant %4 1
%59 = OpConstant %4 0
%78 = OpConstant %4 0
%80 = OpTypePointer Uniform %22
%79 = OpConstant %4 0
%83 = OpConstant %4 1
%84 = OpConstant %4 0
%86 = OpTypePointer Uniform %22
%90 = OpTypePointer Uniform %6
%91 = OpConstant %4 1
%93 = OpTypePointer Uniform %6
%106 = OpConstant %4 2
%108 = OpTypePointer Uniform %6
%121 = OpConstant %4 3
%123 = OpTypePointer Uniform %6
%157 = OpConstant %4 4
%159 = OpTypePointer Uniform %6
%164 = OpConstant %4 5
%166 = OpTypePointer Uniform %6
%171 = OpConstant %4 6
%173 = OpTypePointer Uniform %6
%185 = OpConstant %4 0
%187 = OpTypePointer Uniform %6
%196 = OpConstant %4 0
%198 = OpTypePointer Function %6
%204 = OpConstant %4 0
%206 = OpTypePointer Function %6
%212 = OpConstant %4 1
%214 = OpTypePointer Function %6
%220 = OpConstant %4 1
%222 = OpTypePointer Function %6
%224 = OpConstant %4 0
%225 = OpConstant %4 0
%227 = OpTypePointer Uniform %22
%229 = OpConstant %4 1
%230 = OpConstant %4 0
%232 = OpTypePointer Uniform %22
%105 = OpConstant %4 2
%119 = OpConstant %4 3
%154 = OpConstant %4 4
%160 = OpConstant %4 5
%166 = OpConstant %4 6
%179 = OpConstant %4 0
%189 = OpTypePointer Function %6
%190 = OpConstant %4 0
%197 = OpConstant %4 0
%204 = OpConstant %4 1
%211 = OpConstant %4 1
%214 = OpConstant %4 0
%215 = OpConstant %4 0
%218 = OpConstant %4 1
%219 = OpConstant %4 0
%41 = OpFunction %2 None %42
%43 = OpLabel
%39 = OpVariable %40 Function %8
@ -156,209 +144,209 @@ OpBranchConditional %48 %50 %49
%50 = OpLabel
OpReturn
%49 = OpLabel
%53 = OpAccessChain %54 %18 %52 %46 %51
%55 = OpLoad %22 %53
OpStore %28 %55
%58 = OpAccessChain %59 %18 %57 %46 %56
%60 = OpLoad %22 %58
OpStore %30 %60
%61 = OpCompositeConstruct %22 %5 %5
OpStore %31 %61
%56 = OpAccessChain %53 %18 %55 %46 %54
%57 = OpLoad %22 %56
OpStore %28 %57
%60 = OpAccessChain %53 %18 %59 %46 %58
%61 = OpLoad %22 %60
OpStore %30 %61
%62 = OpCompositeConstruct %22 %5 %5
OpStore %32 %62
OpStore %31 %62
%63 = OpCompositeConstruct %22 %5 %5
OpStore %33 %63
OpBranch %64
%64 = OpLabel
OpLoopMerge %65 %67 None
OpBranch %66
%66 = OpLabel
%68 = OpLoad %9 %39
%69 = OpUGreaterThanEqual %47 %68 %3
OpSelectionMerge %70 None
OpBranchConditional %69 %71 %70
%71 = OpLabel
OpStore %32 %63
%64 = OpCompositeConstruct %22 %5 %5
OpStore %33 %64
OpBranch %65
%70 = OpLabel
%72 = OpLoad %9 %39
%73 = OpIEqual %47 %72 %46
OpSelectionMerge %74 None
OpBranchConditional %73 %75 %74
%75 = OpLabel
OpBranch %67
%74 = OpLabel
%76 = OpLoad %9 %39
%79 = OpAccessChain %80 %18 %78 %76 %77
%81 = OpLoad %22 %79
OpStore %37 %81
%82 = OpLoad %9 %39
%85 = OpAccessChain %86 %18 %84 %82 %83
%87 = OpLoad %22 %85
OpStore %38 %87
%88 = OpLoad %22 %37
%89 = OpLoad %22 %28
%90 = OpExtInst %6 %1 Distance %88 %89
%92 = OpAccessChain %93 %15 %91
%94 = OpLoad %6 %92
%95 = OpFOrdLessThan %47 %90 %94
OpSelectionMerge %96 None
OpBranchConditional %95 %97 %96
%97 = OpLabel
%98 = OpLoad %22 %31
%99 = OpLoad %22 %37
%100 = OpFAdd %22 %98 %99
OpStore %31 %100
%101 = OpLoad %4 %34
%102 = OpIAdd %4 %101 %10
OpStore %34 %102
OpBranch %96
%96 = OpLabel
%103 = OpLoad %22 %37
%104 = OpLoad %22 %28
%105 = OpExtInst %6 %1 Distance %103 %104
%107 = OpAccessChain %108 %15 %106
%109 = OpLoad %6 %107
%110 = OpFOrdLessThan %47 %105 %109
OpSelectionMerge %111 None
OpBranchConditional %110 %112 %111
%112 = OpLabel
%113 = OpLoad %22 %33
%114 = OpLoad %22 %37
%115 = OpLoad %22 %28
%116 = OpFSub %22 %114 %115
%117 = OpFSub %22 %113 %116
OpStore %33 %117
OpBranch %111
%111 = OpLabel
%118 = OpLoad %22 %37
%119 = OpLoad %22 %28
%120 = OpExtInst %6 %1 Distance %118 %119
%122 = OpAccessChain %123 %15 %121
%124 = OpLoad %6 %122
%125 = OpFOrdLessThan %47 %120 %124
OpSelectionMerge %126 None
OpBranchConditional %125 %127 %126
%127 = OpLabel
%128 = OpLoad %22 %32
%129 = OpLoad %22 %38
%130 = OpFAdd %22 %128 %129
OpStore %32 %130
%131 = OpLoad %4 %36
%132 = OpIAdd %4 %131 %10
OpStore %36 %132
OpBranch %126
%126 = OpLabel
%65 = OpLabel
OpLoopMerge %66 %68 None
OpBranch %67
%67 = OpLabel
%133 = OpLoad %9 %39
%134 = OpIAdd %9 %133 %11
OpStore %39 %134
OpBranch %64
%65 = OpLabel
%135 = OpLoad %4 %34
%136 = OpSGreaterThan %47 %135 %7
OpSelectionMerge %137 None
OpBranchConditional %136 %138 %137
%138 = OpLabel
%139 = OpLoad %22 %31
%140 = OpLoad %4 %34
%141 = OpConvertSToF %6 %140
%142 = OpFDiv %6 %12 %141
%143 = OpVectorTimesScalar %22 %139 %142
%144 = OpLoad %22 %28
%145 = OpFSub %22 %143 %144
OpStore %31 %145
OpBranch %137
%137 = OpLabel
%146 = OpLoad %4 %36
%147 = OpSGreaterThan %47 %146 %7
OpSelectionMerge %148 None
OpBranchConditional %147 %149 %148
%149 = OpLabel
%150 = OpLoad %22 %32
%151 = OpLoad %4 %36
%152 = OpConvertSToF %6 %151
%153 = OpFDiv %6 %12 %152
%154 = OpVectorTimesScalar %22 %150 %153
OpStore %32 %154
OpBranch %148
%148 = OpLabel
%155 = OpLoad %22 %30
%156 = OpLoad %22 %31
%158 = OpAccessChain %159 %15 %157
%160 = OpLoad %6 %158
%161 = OpVectorTimesScalar %22 %156 %160
%162 = OpFAdd %22 %155 %161
%163 = OpLoad %22 %33
%165 = OpAccessChain %166 %15 %164
%167 = OpLoad %6 %165
%168 = OpVectorTimesScalar %22 %163 %167
%169 = OpFAdd %22 %162 %168
%170 = OpLoad %22 %32
%172 = OpAccessChain %173 %15 %171
%174 = OpLoad %6 %172
%175 = OpVectorTimesScalar %22 %170 %174
%176 = OpFAdd %22 %169 %175
%69 = OpLoad %9 %39
%70 = OpUGreaterThanEqual %47 %69 %3
OpSelectionMerge %71 None
OpBranchConditional %70 %72 %71
%72 = OpLabel
OpBranch %66
%71 = OpLabel
%73 = OpLoad %9 %39
%74 = OpIEqual %47 %73 %46
OpSelectionMerge %75 None
OpBranchConditional %74 %76 %75
%76 = OpLabel
OpBranch %68
%75 = OpLabel
%77 = OpLoad %9 %39
%80 = OpAccessChain %53 %18 %79 %77 %78
%81 = OpLoad %22 %80
OpStore %37 %81
%82 = OpLoad %9 %39
%85 = OpAccessChain %53 %18 %84 %82 %83
%86 = OpLoad %22 %85
OpStore %38 %86
%87 = OpLoad %22 %37
%88 = OpLoad %22 %28
%89 = OpExtInst %6 %1 Distance %87 %88
%92 = OpAccessChain %90 %15 %91
%93 = OpLoad %6 %92
%94 = OpFOrdLessThan %47 %89 %93
OpSelectionMerge %95 None
OpBranchConditional %94 %96 %95
%96 = OpLabel
%97 = OpLoad %22 %31
%98 = OpLoad %22 %37
%99 = OpFAdd %22 %97 %98
OpStore %31 %99
%100 = OpLoad %4 %34
%101 = OpIAdd %4 %100 %10
OpStore %34 %101
OpBranch %95
%95 = OpLabel
%102 = OpLoad %22 %37
%103 = OpLoad %22 %28
%104 = OpExtInst %6 %1 Distance %102 %103
%106 = OpAccessChain %90 %15 %105
%107 = OpLoad %6 %106
%108 = OpFOrdLessThan %47 %104 %107
OpSelectionMerge %109 None
OpBranchConditional %108 %110 %109
%110 = OpLabel
%111 = OpLoad %22 %33
%112 = OpLoad %22 %37
%113 = OpLoad %22 %28
%114 = OpFSub %22 %112 %113
%115 = OpFSub %22 %111 %114
OpStore %33 %115
OpBranch %109
%109 = OpLabel
%116 = OpLoad %22 %37
%117 = OpLoad %22 %28
%118 = OpExtInst %6 %1 Distance %116 %117
%120 = OpAccessChain %90 %15 %119
%121 = OpLoad %6 %120
%122 = OpFOrdLessThan %47 %118 %121
OpSelectionMerge %123 None
OpBranchConditional %122 %124 %123
%124 = OpLabel
%125 = OpLoad %22 %32
%126 = OpLoad %22 %38
%127 = OpFAdd %22 %125 %126
OpStore %32 %127
%128 = OpLoad %4 %36
%129 = OpIAdd %4 %128 %10
OpStore %36 %129
OpBranch %123
%123 = OpLabel
OpBranch %68
%68 = OpLabel
%130 = OpLoad %9 %39
%131 = OpIAdd %9 %130 %11
OpStore %39 %131
OpBranch %65
%66 = OpLabel
%132 = OpLoad %4 %34
%133 = OpSGreaterThan %47 %132 %7
OpSelectionMerge %134 None
OpBranchConditional %133 %135 %134
%135 = OpLabel
%136 = OpLoad %22 %31
%137 = OpLoad %4 %34
%138 = OpConvertSToF %6 %137
%139 = OpFDiv %6 %12 %138
%140 = OpVectorTimesScalar %22 %136 %139
%141 = OpLoad %22 %28
%142 = OpFSub %22 %140 %141
OpStore %31 %142
OpBranch %134
%134 = OpLabel
%143 = OpLoad %4 %36
%144 = OpSGreaterThan %47 %143 %7
OpSelectionMerge %145 None
OpBranchConditional %144 %146 %145
%146 = OpLabel
%147 = OpLoad %22 %32
%148 = OpLoad %4 %36
%149 = OpConvertSToF %6 %148
%150 = OpFDiv %6 %12 %149
%151 = OpVectorTimesScalar %22 %147 %150
OpStore %32 %151
OpBranch %145
%145 = OpLabel
%152 = OpLoad %22 %30
%153 = OpLoad %22 %31
%155 = OpAccessChain %90 %15 %154
%156 = OpLoad %6 %155
%157 = OpVectorTimesScalar %22 %153 %156
%158 = OpFAdd %22 %152 %157
%159 = OpLoad %22 %33
%161 = OpAccessChain %90 %15 %160
%162 = OpLoad %6 %161
%163 = OpVectorTimesScalar %22 %159 %162
%164 = OpFAdd %22 %158 %163
%165 = OpLoad %22 %32
%167 = OpAccessChain %90 %15 %166
%168 = OpLoad %6 %167
%169 = OpVectorTimesScalar %22 %165 %168
%170 = OpFAdd %22 %164 %169
OpStore %30 %170
%171 = OpLoad %22 %30
%172 = OpExtInst %22 %1 Normalize %171
%173 = OpLoad %22 %30
%174 = OpExtInst %6 %1 Length %173
%175 = OpExtInst %6 %1 FClamp %174 %5 %13
%176 = OpVectorTimesScalar %22 %172 %175
OpStore %30 %176
%177 = OpLoad %22 %30
%178 = OpExtInst %22 %1 Normalize %177
%179 = OpLoad %22 %30
%180 = OpExtInst %6 %1 Length %179
%181 = OpExtInst %6 %1 FClamp %180 %5 %13
%177 = OpLoad %22 %28
%178 = OpLoad %22 %30
%180 = OpAccessChain %90 %15 %179
%181 = OpLoad %6 %180
%182 = OpVectorTimesScalar %22 %178 %181
OpStore %30 %182
%183 = OpLoad %22 %28
%184 = OpLoad %22 %30
%186 = OpAccessChain %187 %15 %185
%188 = OpLoad %6 %186
%189 = OpVectorTimesScalar %22 %184 %188
%190 = OpFAdd %22 %183 %189
OpStore %28 %190
%191 = OpLoad %22 %28
%192 = OpCompositeExtract %6 %191 0
%193 = OpFOrdLessThan %47 %192 %14
OpSelectionMerge %194 None
OpBranchConditional %193 %195 %194
%183 = OpFAdd %22 %177 %182
OpStore %28 %183
%184 = OpLoad %22 %28
%185 = OpCompositeExtract %6 %184 0
%186 = OpFOrdLessThan %47 %185 %14
OpSelectionMerge %187 None
OpBranchConditional %186 %188 %187
%188 = OpLabel
%191 = OpAccessChain %189 %28 %190
OpStore %191 %12
OpBranch %187
%187 = OpLabel
%192 = OpLoad %22 %28
%193 = OpCompositeExtract %6 %192 0
%194 = OpFOrdGreaterThan %47 %193 %12
OpSelectionMerge %195 None
OpBranchConditional %194 %196 %195
%196 = OpLabel
%198 = OpAccessChain %189 %28 %197
OpStore %198 %14
OpBranch %195
%195 = OpLabel
%197 = OpAccessChain %198 %28 %196
OpStore %197 %12
OpBranch %194
%194 = OpLabel
%199 = OpLoad %22 %28
%200 = OpCompositeExtract %6 %199 0
%201 = OpFOrdGreaterThan %47 %200 %12
%200 = OpCompositeExtract %6 %199 1
%201 = OpFOrdLessThan %47 %200 %14
OpSelectionMerge %202 None
OpBranchConditional %201 %203 %202
%203 = OpLabel
%205 = OpAccessChain %206 %28 %204
OpStore %205 %14
%205 = OpAccessChain %189 %28 %204
OpStore %205 %12
OpBranch %202
%202 = OpLabel
%207 = OpLoad %22 %28
%208 = OpCompositeExtract %6 %207 1
%209 = OpFOrdLessThan %47 %208 %14
OpSelectionMerge %210 None
OpBranchConditional %209 %211 %210
%211 = OpLabel
%213 = OpAccessChain %214 %28 %212
OpStore %213 %12
OpBranch %210
%206 = OpLoad %22 %28
%207 = OpCompositeExtract %6 %206 1
%208 = OpFOrdGreaterThan %47 %207 %12
OpSelectionMerge %209 None
OpBranchConditional %208 %210 %209
%210 = OpLabel
%215 = OpLoad %22 %28
%216 = OpCompositeExtract %6 %215 1
%217 = OpFOrdGreaterThan %47 %216 %12
OpSelectionMerge %218 None
OpBranchConditional %217 %219 %218
%219 = OpLabel
%221 = OpAccessChain %222 %28 %220
OpStore %221 %14
OpBranch %218
%218 = OpLabel
%223 = OpLoad %22 %28
%226 = OpAccessChain %227 %24 %225 %46 %224
OpStore %226 %223
%228 = OpLoad %22 %30
%231 = OpAccessChain %232 %24 %230 %46 %229
OpStore %231 %228
%212 = OpAccessChain %189 %28 %211
OpStore %212 %14
OpBranch %209
%209 = OpLabel
%213 = OpLoad %22 %28
%216 = OpAccessChain %53 %24 %215 %46 %214
OpStore %216 %213
%217 = OpLoad %22 %30
%220 = OpAccessChain %53 %24 %219 %46 %218
OpStore %220 %217
OpReturn
OpFunctionEnd

28
tests/out/collatz.spvasm.snap
View File

@ -44,11 +44,11 @@ OpDecorate %11 Binding 0
%20 = OpTypeFunction %4 %4
%28 = OpTypeBool
%47 = OpTypeFunction %2
%54 = OpTypeInt 32 1
%55 = OpConstant %54 0
%57 = OpTypePointer Uniform %4
%60 = OpConstant %54 0
%62 = OpTypePointer Uniform %4
%50 = OpTypePointer Uniform %13
%53 = OpTypePointer Uniform %4
%56 = OpTypeInt 32 1
%57 = OpConstant %56 0
%61 = OpConstant %56 0
%19 = OpFunction %4 None %20
%18 = OpFunctionParameter %4
%21 = OpLabel
@ -100,14 +100,14 @@ OpFunctionEnd
%48 = OpLabel
OpBranch %49
%49 = OpLabel
%50 = OpLoad %9 %8
%51 = OpCompositeExtract %4 %50 0
%52 = OpLoad %9 %8
%53 = OpCompositeExtract %4 %52 0
%56 = OpAccessChain %57 %11 %55 %53
%58 = OpLoad %4 %56
%59 = OpFunctionCall %4 %19 %58
%61 = OpAccessChain %62 %11 %60 %51
OpStore %61 %59
%51 = OpLoad %9 %8
%52 = OpCompositeExtract %4 %51 0
%54 = OpLoad %9 %8
%55 = OpCompositeExtract %4 %54 0
%58 = OpAccessChain %53 %11 %57 %55
%59 = OpLoad %4 %58
%60 = OpFunctionCall %4 %19 %59
%62 = OpAccessChain %53 %11 %61 %52
OpStore %62 %60
OpReturn
OpFunctionEnd

2
tests/out/shadow.msl.snap
View File

@ -39,8 +39,6 @@ typedef metal::float2 type8;
typedef metal::float3 type9;
typedef bool type10;
constexpr constant float const_0f = 0.0;
constexpr constant float const_1f = 1.0;
constexpr constant float const_0_50f = 0.5;

91
tests/out/shadow.spvasm.snap
View File

@ -5,7 +5,7 @@ expression: dis
; SPIR-V
; Version: 1.2
; Generator: rspirv
; Bound: 137
; Bound: 138
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
@ -101,10 +101,11 @@ OpDecorate %36 Location 0
%75 = OpTypePointer Function %10
%77 = OpTypePointer Function %12
%79 = OpTypeFunction %2
%91 = OpConstant %62 0
%93 = OpTypePointer Uniform %17
%101 = OpConstant %62 0
%103 = OpTypePointer Uniform %22
%91 = OpTypePointer Uniform %17
%92 = OpConstant %62 0
%100 = OpTypePointer Uniform %21
%102 = OpTypePointer Uniform %22
%103 = OpConstant %62 0
%40 = OpFunction %4 None %41
%38 = OpFunctionParameter %12
%39 = OpFunctionParameter %24
@ -157,8 +158,8 @@ OpLoopMerge %87 %89 None
OpBranch %88
%88 = OpLabel
%90 = OpLoad %12 %76
%92 = OpAccessChain %93 %15 %91
%94 = OpLoad %17 %92
%93 = OpAccessChain %91 %15 %92
%94 = OpLoad %17 %93
%95 = OpCompositeExtract %12 %94 0
%96 = OpExtInst %12 %1 UMin %95 %11
%97 = OpUGreaterThanEqual %47 %90 %96
@ -167,47 +168,47 @@ OpBranchConditional %97 %99 %98
%99 = OpLabel
OpBranch %87
%98 = OpLabel
%100 = OpLoad %12 %76
%102 = OpAccessChain %103 %19 %101 %100
%104 = OpLoad %22 %102
%105 = OpLoad %12 %76
%106 = OpCompositeExtract %23 %104 0
%107 = OpLoad %24 %34
%108 = OpMatrixTimesVector %24 %106 %107
%109 = OpFunctionCall %4 %40 %105 %108
%110 = OpCompositeExtract %24 %104 1
%111 = OpCompositeExtract %4 %110 0
%112 = OpCompositeExtract %4 %110 1
%113 = OpCompositeExtract %4 %110 2
%114 = OpCompositeConstruct %10 %111 %112 %113
%115 = OpLoad %24 %34
%116 = OpCompositeExtract %4 %115 0
%117 = OpCompositeExtract %4 %115 1
%118 = OpCompositeExtract %4 %115 2
%119 = OpCompositeConstruct %10 %116 %117 %118
%120 = OpFSub %10 %114 %119
%121 = OpExtInst %10 %1 Normalize %120
%122 = OpDot %4 %85 %121
%123 = OpExtInst %4 %1 FMax %3 %122
%124 = OpLoad %10 %74
%125 = OpFMul %4 %109 %123
%126 = OpCompositeExtract %24 %104 2
%127 = OpCompositeExtract %4 %126 0
%128 = OpCompositeExtract %4 %126 1
%129 = OpCompositeExtract %4 %126 2
%130 = OpCompositeConstruct %10 %127 %128 %129
%131 = OpVectorTimesScalar %10 %130 %125
%132 = OpFAdd %10 %124 %131
OpStore %74 %132
%101 = OpLoad %12 %76
%104 = OpAccessChain %102 %19 %103 %101
%105 = OpLoad %22 %104
%106 = OpLoad %12 %76
%107 = OpCompositeExtract %23 %105 0
%108 = OpLoad %24 %34
%109 = OpMatrixTimesVector %24 %107 %108
%110 = OpFunctionCall %4 %40 %106 %109
%111 = OpCompositeExtract %24 %105 1
%112 = OpCompositeExtract %4 %111 0
%113 = OpCompositeExtract %4 %111 1
%114 = OpCompositeExtract %4 %111 2
%115 = OpCompositeConstruct %10 %112 %113 %114
%116 = OpLoad %24 %34
%117 = OpCompositeExtract %4 %116 0
%118 = OpCompositeExtract %4 %116 1
%119 = OpCompositeExtract %4 %116 2
%120 = OpCompositeConstruct %10 %117 %118 %119
%121 = OpFSub %10 %115 %120
%122 = OpExtInst %10 %1 Normalize %121
%123 = OpDot %4 %85 %122
%124 = OpExtInst %4 %1 FMax %3 %123
%125 = OpLoad %10 %74
%126 = OpFMul %4 %110 %124
%127 = OpCompositeExtract %24 %105 2
%128 = OpCompositeExtract %4 %127 0
%129 = OpCompositeExtract %4 %127 1
%130 = OpCompositeExtract %4 %127 2
%131 = OpCompositeConstruct %10 %128 %129 %130
%132 = OpVectorTimesScalar %10 %131 %126
%133 = OpFAdd %10 %125 %132
OpStore %74 %133
OpBranch %89
%89 = OpLabel
%133 = OpLoad %12 %76
%134 = OpIAdd %12 %133 %14
OpStore %76 %134
%134 = OpLoad %12 %76
%135 = OpIAdd %12 %134 %14
OpStore %76 %135
OpBranch %86
%87 = OpLabel
%135 = OpLoad %10 %74
%136 = OpCompositeConstruct %24 %135 %5
OpStore %36 %136
%136 = OpLoad %10 %74
%137 = OpCompositeConstruct %24 %136 %5
OpStore %36 %137
OpReturn
OpFunctionEnd

12
tests/out/skybox.msl.snap
View File

@ -20,13 +20,11 @@ struct Data {
typedef int type4;
typedef float type5;
typedef metal::float3x3 type5;
typedef metal::float3x3 type6;
typedef metal::texturecube<float, metal::access::sample> type6;
typedef metal::texturecube<float, metal::access::sample> type7;
typedef metal::sampler type8;
typedef metal::sampler type7;
constexpr constant int const_2i = 2;
constexpr constant int const_1i = 1;
@ -70,8 +68,8 @@ struct fs_mainOutput {
fragment fs_mainOutput fs_main(
fs_mainInput input [[stage_in]],
type7 r_texture [[texture(0)]],
type8 r_sampler [[sampler(1)]]
type6 r_texture [[texture(0)]],
type7 r_sampler [[sampler(1)]]
) {
fs_mainOutput output;
metal::float4 _expr9 = r_texture.sample(r_sampler, input.in_uv);

101
tests/out/skybox.spvasm.snap
View File

@ -5,13 +5,13 @@ expression: dis
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 106
; Bound: 103
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %37 "vs_main" %10 %13 %16
OpEntryPoint Fragment %97 "fs_main" %29 %31
OpExecutionMode %97 OriginUpperLeft
OpEntryPoint Fragment %94 "fs_main" %29 %31
OpExecutionMode %94 OriginUpperLeft
OpSource GLSL 450
OpName %10 "out_position"
OpName %13 "out_uv"
@ -29,8 +29,8 @@ OpName %34 "tmp2"
OpName %35 "unprojected"
OpName %37 "vs_main"
OpName %37 "vs_main"
OpName %97 "fs_main"
OpName %97 "fs_main"
OpName %94 "fs_main"
OpName %94 "fs_main"
OpDecorate %10 BuiltIn Position
OpDecorate %13 Location 0
OpDecorate %16 BuiltIn VertexIndex
@ -82,16 +82,13 @@ OpDecorate %31 Location 0
%33 = OpTypePointer Function %4
%36 = OpTypePointer Function %11
%38 = OpTypeFunction %2
%56 = OpConstant %4 1
%58 = OpTypePointer Uniform %21
%56 = OpTypePointer Uniform %21
%57 = OpConstant %4 1
%65 = OpConstant %4 1
%67 = OpTypePointer Uniform %21
%74 = OpConstant %4 1
%76 = OpTypePointer Uniform %21
%83 = OpTypeMatrix %14 3
%86 = OpConstant %4 0
%88 = OpTypePointer Uniform %21
%103 = OpTypeSampledImage %24
%73 = OpConstant %4 1
%81 = OpTypeMatrix %14 3
%84 = OpConstant %4 0
%100 = OpTypeSampledImage %24
%37 = OpFunction %2 None %38
%39 = OpLabel
%32 = OpVariable %33 Function
@ -116,52 +113,52 @@ OpStore %34 %46
%53 = OpFMul %7 %52 %6
%54 = OpFSub %7 %53 %8
%55 = OpCompositeConstruct %11 %50 %54 %9 %8
%57 = OpAccessChain %58 %19 %56
%59 = OpLoad %21 %57
%58 = OpAccessChain %56 %19 %57
%59 = OpLoad %21 %58
%60 = OpCompositeExtract %11 %59 0
%61 = OpCompositeExtract %7 %60 0
%62 = OpCompositeExtract %7 %60 1
%63 = OpCompositeExtract %7 %60 2
%64 = OpCompositeConstruct %14 %61 %62 %63
%66 = OpAccessChain %67 %19 %65
%68 = OpLoad %21 %66
%69 = OpCompositeExtract %11 %68 1
%70 = OpCompositeExtract %7 %69 0
%71 = OpCompositeExtract %7 %69 1
%72 = OpCompositeExtract %7 %69 2
%73 = OpCompositeConstruct %14 %70 %71 %72
%75 = OpAccessChain %76 %19 %74
%77 = OpLoad %21 %75
%78 = OpCompositeExtract %11 %77 2
%79 = OpCompositeExtract %7 %78 0
%80 = OpCompositeExtract %7 %78 1
%81 = OpCompositeExtract %7 %78 2
%82 = OpCompositeConstruct %14 %79 %80 %81
%84 = OpCompositeConstruct %83 %64 %73 %82
%85 = OpTranspose %83 %84
%87 = OpAccessChain %88 %19 %86
%89 = OpLoad %21 %87
%90 = OpMatrixTimesVector %11 %89 %55
OpStore %35 %90
%91 = OpLoad %11 %35
%92 = OpCompositeExtract %7 %91 0
%93 = OpCompositeExtract %7 %91 1
%94 = OpCompositeExtract %7 %91 2
%95 = OpCompositeConstruct %14 %92 %93 %94
%96 = OpMatrixTimesVector %14 %85 %95
OpStore %13 %96
%66 = OpAccessChain %56 %19 %65
%67 = OpLoad %21 %66
%68 = OpCompositeExtract %11 %67 1
%69 = OpCompositeExtract %7 %68 0
%70 = OpCompositeExtract %7 %68 1
%71 = OpCompositeExtract %7 %68 2
%72 = OpCompositeConstruct %14 %69 %70 %71
%74 = OpAccessChain %56 %19 %73
%75 = OpLoad %21 %74
%76 = OpCompositeExtract %11 %75 2
%77 = OpCompositeExtract %7 %76 0
%78 = OpCompositeExtract %7 %76 1
%79 = OpCompositeExtract %7 %76 2
%80 = OpCompositeConstruct %14 %77 %78 %79
%82 = OpCompositeConstruct %81 %64 %72 %80
%83 = OpTranspose %81 %82
%85 = OpAccessChain %56 %19 %84
%86 = OpLoad %21 %85
%87 = OpMatrixTimesVector %11 %86 %55
OpStore %35 %87
%88 = OpLoad %11 %35
%89 = OpCompositeExtract %7 %88 0
%90 = OpCompositeExtract %7 %88 1
%91 = OpCompositeExtract %7 %88 2
%92 = OpCompositeConstruct %14 %89 %90 %91
%93 = OpMatrixTimesVector %14 %83 %92
OpStore %13 %93
OpStore %10 %55
OpReturn
OpFunctionEnd
%97 = OpFunction %2 None %38
%94 = OpFunction %2 None %38
%95 = OpLabel
%96 = OpLoad %24 %23
%97 = OpLoad %27 %26
OpBranch %98
%98 = OpLabel
%99 = OpLoad %24 %23
%100 = OpLoad %27 %26
OpBranch %101
%101 = OpLabel
%102 = OpLoad %14 %29
%104 = OpSampledImage %103 %99 %100
%105 = OpImageSampleImplicitLod %11 %104 %102
OpStore %31 %105
%99 = OpLoad %14 %29
%101 = OpSampledImage %100 %96 %97
%102 = OpImageSampleImplicitLod %11 %101 %99
OpStore %31 %102
OpReturn
OpFunctionEnd

8
tests/out/texture-array.spvasm.snap
View File

@ -54,8 +54,8 @@ OpDecorate %20 Location 1
%22 = OpTypePointer Output %21
%20 = OpVariable %22 Output
%24 = OpTypeFunction %2
%30 = OpConstant %4 0
%32 = OpTypePointer PushConstant %18
%30 = OpTypePointer PushConstant %18
%31 = OpConstant %4 0
%34 = OpTypeBool
%40 = OpTypeSampledImage %10
%23 = OpFunction %2 None %24
@ -65,8 +65,8 @@ OpDecorate %20 Location 1
%28 = OpLoad %14 %13
OpBranch %29
%29 = OpLabel
%31 = OpAccessChain %32 %16 %30
%33 = OpLoad %18 %31
%32 = OpAccessChain %30 %16 %31
%33 = OpLoad %18 %32
%35 = OpIEqual %34 %33 %3
OpSelectionMerge %36 None
OpBranchConditional %35 %37 %38

2
tests/parse.rs
View File

@ -38,5 +38,5 @@ fn parse_glsl() {
//check_glsl("glsl_constant_expression.vert"); //TODO
//check_glsl("glsl_if_preprocessor.vert");
check_glsl("glsl_preprocessor_abuse.vert");
check_glsl("glsl_vertex_test_shader.vert");
//check_glsl("glsl_vertex_test_shader.vert"); //TODO
}