nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-10-30 14:01:39 +00:00
Code Issues Packages Projects Releases Wiki Activity

Remove top_level property of structs

Browse Source
This commit is contained in:
Dzmitry Malyshau 2021-12-08 00:27:20 -05:00
parent 3867ef4f6c
commit 8ffd6ba929
78 changed files with 888 additions and 986 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
CHANGELOG.md
View File

@ -1,6 +1,10 @@
# Change Log
## v0.7.1 (2021-10-12)
## v0.8 (TBD)
- WGSL-in:
- remove `[[block]]` attribute
### v0.7.1 (2021-10-12)
- implement casts from and to booleans in the backends
## v0.7 (2021-10-07)

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

@ -83,6 +83,33 @@ impl crate::AtomicFunction {
}
}
impl crate::StorageClass {
fn is_buffer(&self) -> bool {
match *self {
crate::StorageClass::Uniform | crate::StorageClass::Storage { .. } => true,
_ => false,
}
}
}
//Note: similar to `back/spv/helpers.rs`
fn global_needs_wrapper(ir_module: &crate::Module, global_ty: Handle<crate::Type>) -> bool {
match ir_module.types[global_ty].inner {
crate::TypeInner::Struct {
ref members,
span: _,
} => match ir_module.types[members.last().unwrap().ty].inner {
// Structs with dynamically sized arrays can't be copied and can't be wrapped.
crate::TypeInner::Array {
size: crate::ArraySize::Dynamic,
..
} => false,
_ => true,
},
_ => false,
}
}
/// glsl version
#[derive(Debug, Copy, Clone, PartialEq)]
#[cfg_attr(feature = "serialize", derive(serde::Serialize))]
@ -340,8 +367,6 @@ pub struct Writer<'a, W> {
/// A map with all the names needed for writing the module
/// (generated by a [`Namer`](crate::proc::Namer))
names: crate::FastHashMap<NameKey, String>,
/// A map with all the names needed for reflections
reflection_names_uniforms: crate::FastHashMap<Handle<crate::Type>, String>,
/// A map with the names of global variables needed for reflections
reflection_names_globals: crate::FastHashMap<Handle<crate::GlobalVariable>, String>,
/// The selected entry point
@ -397,7 +422,6 @@ impl<'a, W: Write> Writer<'a, W> {
namer,
features: FeaturesManager::new(),
names,
reflection_names_uniforms: crate::FastHashMap::default(),
reflection_names_globals: crate::FastHashMap::default(),
entry_point: &module.entry_points[ep_idx],
entry_point_idx: ep_idx as u16,
@ -489,27 +513,30 @@ impl<'a, W: Write> Writer<'a, W> {
}
}
let ep_info = self.info.get_entry_point(self.entry_point_idx as usize);
// Write all structs
//
// This are always ordered because of the IR is structured in a way that you can't make a
// struct without adding all of it's members first
for (handle, ty) in self.module.types.iter() {
if let TypeInner::Struct { ref members, .. } = ty.inner {
// No needed to write a struct that also should be written as a global variable
let is_global_struct = self
.module
.global_variables
.iter()
.any(|e| e.1.ty == handle);
let used_by_global = self.module.global_variables.iter().any(|(vh, var)| {
!ep_info[vh].is_empty() && var.class.is_buffer() && var.ty == handle
});
if !is_global_struct {
self.write_struct(false, handle, members)?
let is_wrapped = global_needs_wrapper(self.module, handle);
// If it's a global non-wrapped struct, it will be printed
// with the corresponding global variable.
if !used_by_global || is_wrapped {
let name = &self.names[&NameKey::Type(handle)];
write!(self.out, "struct {} ", name)?;
self.write_struct_body(handle, members)?;
writeln!(self.out, ";")?;
}
}
}
let ep_info = self.info.get_entry_point(self.entry_point_idx as usize);
// Write the globals
//
// We filter all globals that aren't used by the selected entry point as they might be
@ -748,12 +775,7 @@ impl<'a, W: Write> Writer<'a, W> {
match self.module.types[ty].inner {
// glsl has no pointer types so just write types as normal and loads are skipped
TypeInner::Pointer { base, .. } => self.write_type(base),
TypeInner::Struct {
top_level: true,
ref members,
span: _,
} => self.write_struct(true, ty, members),
// glsl structs are written as just the struct name if it isn't a block
// glsl structs are written as just the struct name
TypeInner::Struct { .. } => {
// Get the struct name
let name = &self.names[&NameKey::Type(ty)];
@ -861,21 +883,46 @@ impl<'a, W: Write> Writer<'a, W> {
// Trailing space is important
if let Some(storage_class) = glsl_storage_class(global.class) {
write!(self.out, "{} ", storage_class)?;
} else if let TypeInner::Struct {
top_level: true, ..
} = self.module.types[global.ty].inner
{
write!(self.out, "struct ")?;
}
// Write the type
// `write_type` adds no leading or trailing spaces
self.write_type(global.ty)?;
// If struct is a block we need to write `block_name { members }` where `block_name` must be
// unique between blocks and structs so we add `_block_ID` where `ID` is a `IdGenerator`
// generated number so it's unique and `members` are the same as in a struct
// Write the block name, it's just the struct name appended with `_block_ID`
let needs_wrapper = if global.class.is_buffer() {
let ty_name = &self.names[&NameKey::Type(global.ty)];
let block_name = format!(
"{}_block_{}{:?}",
ty_name,
self.block_id.generate(),
self.entry_point.stage,
);
write!(self.out, "{} ", block_name)?;
self.reflection_names_globals.insert(handle, block_name);
let needs_wrapper = global_needs_wrapper(self.module, global.ty);
if needs_wrapper {
write!(self.out, "{{ ")?;
// Write the type
// `write_type` adds no leading or trailing spaces
self.write_type(global.ty)?;
} else if let crate::TypeInner::Struct { ref members, .. } =
self.module.types[global.ty].inner
{
self.write_struct_body(global.ty, members)?;
}
needs_wrapper
} else {
self.write_type(global.ty)?;
false
};
// Finally write the global name and end the global with a `;` and a newline
// Leading space is important
write!(self.out, " ")?;
self.write_global_name(handle, global)?;
if let TypeInner::Array { size, .. } = self.module.types[global.ty].inner {
self.write_array_size(size)?;
}
@ -888,8 +935,12 @@ impl<'a, W: Write> Writer<'a, W> {
self.write_zero_init_value(global.ty)?;
}
}
writeln!(self.out, ";")?;
if needs_wrapper {
write!(self.out, "; }}")?;
}
writeln!(self.out, ";")?;
Ok(())
}
@ -1283,9 +1334,8 @@ impl<'a, W: Write> Writer<'a, W> {
///
/// # Notes
/// Ends in a newline
fn write_struct(
fn write_struct_body(
&mut self,
block: bool,
handle: Handle<crate::Type>,
members: &[crate::StructMember],
) -> BackendResult {
@ -1296,30 +1346,7 @@ impl<'a, W: Write> Writer<'a, W> {
// | `members` is a semicolon separated list of `type name`
// | `type` is the member type
// | `name` is the member name
let name = &self.names[&NameKey::Type(handle)];
// If struct is a block we need to write `block_name { members }` where `block_name` must be
// unique between blocks and structs so we add `_block_ID` where `ID` is a `IdGenerator`
// generated number so it's unique and `members` are the same as in a struct
if block {
// Write the block name, it's just the struct name appended with `_block_ID`
let stage_postfix = match self.entry_point.stage {
ShaderStage::Vertex => "Vs",
ShaderStage::Fragment => "Fs",
ShaderStage::Compute => "Cs",
};
let block_name = format!(
"{}_block_{}{}",
name,
self.block_id.generate(),
stage_postfix
);
writeln!(self.out, "{} {{", block_name)?;
self.reflection_names_uniforms.insert(handle, block_name);
} else {
writeln!(self.out, "struct {} {{", name)?;
}
writeln!(self.out, "{{")?;
for (idx, member) in members.iter().enumerate() {
// The indentation is only for readability
@ -1360,13 +1387,6 @@ impl<'a, W: Write> Writer<'a, W> {
}
write!(self.out, "}}")?;
if !block {
writeln!(self.out, ";")?;
// Add a newline for readability
writeln!(self.out)?;
}
Ok(())
}
@ -2747,7 +2767,7 @@ impl<'a, W: Write> Writer<'a, W> {
match self.module.types[var.ty].inner {
crate::TypeInner::Struct { .. } => match var.class {
crate::StorageClass::Uniform | crate::StorageClass::Storage { .. } => {
let name = self.reflection_names_uniforms[&var.ty].clone();
let name = self.reflection_names_globals[&handle].clone();
uniforms.insert(handle, name);
}
_ => (),

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

@ -152,12 +152,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
// Write all structs
for (handle, ty) in module.types.iter() {
if let TypeInner::Struct {
top_level,
ref members,
..
} = ty.inner
{
if let TypeInner::Struct { ref members, .. } = ty.inner {
if let Some(member) = members.last() {
if let TypeInner::Array {
size: crate::ArraySize::Dynamic,
@ -180,7 +175,6 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
self.write_struct(
module,
handle,
top_level,
members,
ep_result.map(|r| (r.0, Io::Output)),
)?;
@ -703,7 +697,6 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
&mut self,
module: &Module,
handle: Handle<crate::Type>,
_block: bool,
members: &[crate::StructMember],
shader_stage: Option<(ShaderStage, Io)>,
) -> BackendResult {
@ -1941,11 +1934,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
let var = &module.global_variables[var_handle];
let (offset, stride) = match module.types[var.ty].inner {
TypeInner::Array { stride, .. } => (0, stride),
TypeInner::Struct {
top_level: true,
ref members,
..
} => {
TypeInner::Struct { ref members, .. } => {
let last = members.last().unwrap();
let stride = match module.types[last.ty].inner {
TypeInner::Array { stride, .. } => stride,

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

@ -2652,9 +2652,7 @@ impl<W: Write> Writer<W> {
if let Some(ref br) = var.binding {
let good = match options.per_stage_map[ep.stage].resources.get(br) {
Some(target) => match module.types[var.ty].inner {
crate::TypeInner::Struct {
top_level: true, ..
} => target.buffer.is_some(),
crate::TypeInner::Struct { .. } => target.buffer.is_some(),
crate::TypeInner::Image { .. } => target.texture.is_some(),
crate::TypeInner::Sampler { .. } => target.sampler.is_some(),
_ => false,

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

@ -247,7 +247,7 @@ impl<'w> BlockContext<'w> {
}
}
crate::Expression::GlobalVariable(handle) => {
self.writer.global_variables[handle.index()].id
self.writer.global_variables[handle.index()].access_id
}
crate::Expression::Constant(handle) => self.writer.constant_ids[handle.index()],
crate::Expression::Splat { size, value } => {
@ -1065,7 +1065,7 @@ impl<'w> BlockContext<'w> {
}
crate::Expression::GlobalVariable(handle) => {
let gv = &self.writer.global_variables[handle.index()];
break gv.id;
break gv.access_id;
}
crate::Expression::LocalVariable(variable) => {
let local_var = &self.function.variables[&variable];

26
src/back/spv/helpers.rs
View File

@ -61,3 +61,29 @@ impl crate::StorageClass {
}
}
}
/// Return true if the global requires a type decorated with "Block".
// See `back::spv::GlobalVariable::access_id` for details.
pub fn global_needs_wrapper(ir_module: &crate::Module, var: &crate::GlobalVariable) -> bool {
match var.class {
crate::StorageClass::Uniform | crate::StorageClass::Storage { .. } => {}
_ => return false,
};
match ir_module.types[var.ty].inner {
crate::TypeInner::Struct {
ref members,
span: _,
} => match members.last() {
Some(member) => match ir_module.types[member.ty].inner {
// Structs with dynamically sized arrays can't be copied and can't be wrapped.
crate::TypeInner::Array {
size: crate::ArraySize::Dynamic,
..
} => false,
_ => true,
},
None => false,
},
_ => false,
}
}

7
src/back/spv/index.rs
View File

@ -43,9 +43,10 @@ impl<'w> BlockContext<'w> {
let (structure_id, last_member_index) = match self.ir_function.expressions[array] {
crate::Expression::AccessIndex { base, index } => {
match self.ir_function.expressions[base] {
crate::Expression::GlobalVariable(handle) => {
(self.writer.global_variables[handle.index()].id, index)
}
crate::Expression::GlobalVariable(handle) => (
self.writer.global_variables[handle.index()].access_id,
index,
),
_ => return Err(Error::Validation("array length expression")),
}
}

27
src/back/spv/mod.rs
View File

@ -430,30 +430,49 @@ impl recyclable::Recyclable for CachedExpressions {
}
}
#[derive(Clone)]
struct GlobalVariable {
/// Actual ID of the variable.
id: Word,
/// ID of the variable. Not really used.
var_id: Word,
/// For `StorageClass::Handle` variables, this ID is recorded in the function
/// prelude block (and reset before every function) as `OpLoad` of the variable.
/// It is then used for all the global ops, such as `OpImageSample`.
handle_id: Word,
/// Actual ID used to access this variable.
/// For wrapped buffer variables, this ID is `OpAccessChain` into the
/// wrapper. Otherwise, the same as `var_id`.
///
/// Vulkan requires that globals in the `StorageBuffer` and `Uniform` storage
/// classes must be structs with the `Block` decoration, but WGSL and Naga IR
/// make no such requirement. So for such variables, we generate a wrapper struct
/// type with a single element of the type given by Naga, generate an
/// `OpAccessChain` for that member in the function prelude, and use that pointer
/// to refer to the global in the function body. This is the id of that access,
/// updated for each function in `write_function`.
access_id: Word,
}
impl GlobalVariable {
fn dummy() -> Self {
Self {
id: 0,
var_id: 0,
handle_id: 0,
access_id: 0,
}
}
fn new(id: Word) -> Self {
Self { id, handle_id: 0 }
Self {
var_id: id,
handle_id: 0,
access_id: 0,
}
}
/// Prepare `self` for use within a single function.
fn reset_for_function(&mut self) {
self.handle_id = 0;
self.access_id = 0;
}
}

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

@ -1,5 +1,5 @@
use super::{
helpers::{contains_builtin, map_storage_class},
helpers::{contains_builtin, global_needs_wrapper, map_storage_class},
make_local, Block, BlockContext, CachedExpressions, EntryPointContext, Error, Function,
FunctionArgument, GlobalVariable, IdGenerator, Instruction, LocalType, LocalVariable,
LogicalLayout, LookupFunctionType, LookupType, LoopContext, Options, PhysicalLayout,
@ -485,22 +485,46 @@ impl Writer {
function.entry_point_context = Some(ep_context);
}
// fill up the `GlobalVariable::handle_id`
// fill up the `GlobalVariable::access_id`
for gv in self.global_variables.iter_mut() {
gv.reset_for_function();
}
for (handle, var) in ir_module.global_variables.iter() {
// Handle globals are pre-emitted and should be loaded automatically.
if info[handle].is_empty() || var.class != crate::StorageClass::Handle {
if info[handle].is_empty() {
continue;
}
let id = self.id_gen.next();
let result_type_id = self.get_type_id(LookupType::Handle(var.ty));
let gv = &mut self.global_variables[handle.index()];
prelude
.body
.push(Instruction::load(result_type_id, id, gv.id, None));
gv.handle_id = id;
let mut gv = self.global_variables[handle.index()].clone();
// Handle globals are pre-emitted and should be loaded automatically.
if var.class == crate::StorageClass::Handle {
let var_type_id = self.get_type_id(LookupType::Handle(var.ty));
let id = self.id_gen.next();
prelude
.body
.push(Instruction::load(var_type_id, id, gv.var_id, None));
gv.access_id = gv.var_id;
gv.handle_id = id;
} else if global_needs_wrapper(ir_module, var) {
let class = map_storage_class(var.class);
let pointer_type_id = self.get_pointer_id(&ir_module.types, var.ty, class)?;
let index_id = self.get_index_constant(0);
let id = self.id_gen.next();
prelude.body.push(Instruction::access_chain(
pointer_type_id,
id,
gv.var_id,
&[index_id],
));
gv.access_id = id;
} else {
// by default, the variable ID is accessed as is
gv.access_id = gv.var_id;
};
// work around borrow checking in the presense of `self.xxx()` calls
self.global_variables[handle.index()] = gv;
}
// Create a `BlockContext` for generating SPIR-V for the function's
@ -818,14 +842,9 @@ impl Writer {
}
}
crate::TypeInner::Struct {
top_level,
ref members,
span: _,
} => {
if top_level {
self.decorate(id, Decoration::Block, &[]);
}
let mut member_ids = Vec::with_capacity(members.len());
for (index, member) in members.iter().enumerate() {
if decorate_layout {
@ -1215,17 +1234,13 @@ impl Writer {
&mut self,
ir_module: &crate::Module,
global_variable: &crate::GlobalVariable,
) -> Result<(Instruction, Word), Error> {
) -> Result<Word, Error> {
use spirv::Decoration;
let id = self.id_gen.next();
let class = map_storage_class(global_variable.class);
//self.check(class.required_capabilities())?;
let init_word = global_variable
.init
.map(|constant| self.constant_ids[constant.index()]);
let pointer_type_id = self.get_pointer_id(&ir_module.types, global_variable.ty, class)?;
let instruction = Instruction::variable(pointer_type_id, id, class, init_word);
//self.check(class.required_capabilities())?;
if self.flags.contains(WriterFlags::DEBUG) {
if let Some(ref name) = global_variable.name {
@ -1233,8 +1248,6 @@ impl Writer {
}
}
use spirv::Decoration;
let storage_access = match global_variable.class {
crate::StorageClass::Storage { access } => Some(access),
_ => match ir_module.types[global_variable.ty].inner {
@ -1259,8 +1272,43 @@ impl Writer {
self.decorate(id, Decoration::Binding, &[res_binding.binding]);
}
// TODO Initializer is optional and not (yet) included in the IR
Ok((instruction, id))
let init_word = global_variable
.init
.map(|constant| self.constant_ids[constant.index()]);
let inner_type_id = self.get_type_id(LookupType::Handle(global_variable.ty));
// generate the wrapping structure if needed
let pointer_type_id = if global_needs_wrapper(ir_module, global_variable) {
let wrapper_type_id = self.id_gen.next();
self.decorate(wrapper_type_id, Decoration::Block, &[]);
self.annotations.push(Instruction::member_decorate(
wrapper_type_id,
0,
Decoration::Offset,
&[0],
));
Instruction::type_struct(wrapper_type_id, &[inner_type_id])
.to_words(&mut self.logical_layout.declarations);
let pointer_type_id = self.id_gen.next();
Instruction::type_pointer(pointer_type_id, class, wrapper_type_id)
.to_words(&mut self.logical_layout.declarations);
pointer_type_id
} else {
// This is a global variable in a Storage class. The only way it could
// have `global_needs_wrapper() == false` is if it has a runtime-sized array.
// In this case, we need to decorate it with Block.
if let crate::StorageClass::Storage { .. } = global_variable.class {
self.decorate(inner_type_id, Decoration::Block, &[]);
}
self.get_pointer_id(&ir_module.types, global_variable.ty, class)?
};
Instruction::variable(pointer_type_id, id, class, init_word)
.to_words(&mut self.logical_layout.declarations);
Ok(id)
}
fn get_function_type(&mut self, lookup_function_type: LookupFunctionType) -> Word {
@ -1389,8 +1437,7 @@ impl Writer {
GlobalVariable::dummy()
}
_ => {
let (instruction, id) = self.write_global_variable(ir_module, var)?;
instruction.to_words(&mut self.logical_layout.declarations);
let id = self.write_global_variable(ir_module, var)?;
GlobalVariable::new(id)
}
};

13
src/back/wgsl/writer.rs
View File

@ -12,7 +12,6 @@ type BackendResult = Result<(), Error>;
/// WGSL [attribute](https://gpuweb.github.io/gpuweb/wgsl/#attributes)
enum Attribute {
Binding(u32),
Block,
BuiltIn(crate::BuiltIn),
Group(u32),
Interpolate(Option<crate::Interpolation>, Option<crate::Sampling>),
@ -107,12 +106,11 @@ impl<W: Write> Writer<W> {
// Write all structs
for (handle, ty) in module.types.iter() {
if let TypeInner::Struct {
top_level,
ref members,
..
span: _,
} = ty.inner
{
self.write_struct(module, handle, top_level, members)?;
self.write_struct(module, handle, members)?;
writeln!(self.out)?;
}
}
@ -358,7 +356,6 @@ impl<W: Write> Writer<W> {
for (index, attribute) in attributes.iter().enumerate() {
match *attribute {
Attribute::Block => write!(self.out, "block")?,
Attribute::Location(id) => write!(self.out, "location({})", id)?,
Attribute::BuiltIn(builtin_attrib) => {
if let Some(builtin) = builtin_str(builtin_attrib) {
@ -432,14 +429,8 @@ impl<W: Write> Writer<W> {
&mut self,
module: &Module,
handle: Handle<crate::Type>,
block: bool,
members: &[crate::StructMember],
) -> BackendResult {
if block {
self.write_attributes(&[Attribute::Block], false)?;
writeln!(self.out)?;
}
write!(self.out, "struct ")?;
self.write_struct_name(module, handle)?;
write!(self.out, " {{")?;

6
src/front/glsl/functions.rs
View File

@ -1141,11 +1141,7 @@ impl Parser {
let ty = self.module.types.insert(
Type {
name: None,
inner: TypeInner::Struct {
top_level: false,
members,
span,
},
inner: TypeInner::Struct { members, span },
},
Default::default(),
);

12
src/front/glsl/offset.rs
View File

@ -120,11 +120,7 @@ pub fn calculate_offset(
(align, align * columns as u32)
}
TypeInner::Struct {
ref members,
top_level,
..
} => {
TypeInner::Struct { ref members, .. } => {
let mut span = 0;
let mut align = 0;
let mut members = members.clone();
@ -148,11 +144,7 @@ pub fn calculate_offset(
ty = types.insert(
Type {
name,
inner: TypeInner::Struct {
top_level,
members,
span,
},
inner: TypeInner::Struct { members, span },
},
ty_span,
);

1
src/front/glsl/parser/declarations.rs
View File

@ -501,7 +501,6 @@ impl<'source> ParsingContext<'source> {
Type {
name: Some(ty_name),
inner: TypeInner::Struct {
top_level: true,
members: members.clone(),
span,
},

6
src/front/glsl/parser/types.rs
View File

@ -61,11 +61,7 @@ impl<'source> ParsingContext<'source> {
let ty = parser.module.types.insert(
Type {
name: Some(ty_name.clone()),
inner: TypeInner::Struct {
top_level: false,
members,
span,
},
inner: TypeInner::Struct { members, span },
},
meta,
);

1
src/front/spv/function.rs
View File

@ -488,7 +488,6 @@ impl<I: Iterator<Item = u32>> super::Parser<I> {
crate::Type {
name: None,
inner: crate::TypeInner::Struct {
top_level: false,
members,
span: 0xFFFF, // shouldn't matter
},

24
src/front/spv/mod.rs
View File

@ -160,11 +160,6 @@ impl crate::ImageDimension {
type MemberIndex = u32;
#[derive(Clone, Debug, Default, PartialEq)]
struct Block {
buffer: bool,
}
bitflags::bitflags! {
#[derive(Default)]
struct DecorationFlags: u32 {
@ -200,7 +195,7 @@ struct Decoration {
desc_set: Option<spirv::Word>,
desc_index: Option<spirv::Word>,
specialization: Option<spirv::Word>,
block: Option<Block>,
storage_buffer: bool,
offset: Option<spirv::Word>,
array_stride: Option<NonZeroU32>,
matrix_stride: Option<NonZeroU32>,
@ -663,11 +658,8 @@ impl<I: Iterator<Item = u32>> Parser<I> {
inst.expect(base_words + 2)?;
dec.desc_index = Some(self.next()?);
}
spirv::Decoration::Block => {
dec.block = Some(Block { buffer: false });
}
spirv::Decoration::BufferBlock => {
dec.block = Some(Block { buffer: true });
dec.storage_buffer = true;
}
spirv::Decoration::Offset => {
inst.expect(base_words + 2)?;
@ -4022,7 +4014,9 @@ impl<I: Iterator<Item = u32>> Parser<I> {
inst.expect_at_least(2)?;
let id = self.next()?;
let parent_decor = self.future_decor.remove(&id);
let block_decor = parent_decor.as_ref().and_then(|decor| decor.block.clone());
let is_storage_buffer = parent_decor
.as_ref()
.map_or(false, |decor| decor.storage_buffer);
self.layouter
.update(&module.types, &module.constants)
@ -4097,11 +4091,7 @@ impl<I: Iterator<Item = u32>> Parser<I> {
span = crate::front::align_up(span, alignment);
let inner = crate::TypeInner::Struct {
top_level: block_decor.is_some(),
span,
members,
};
let inner = crate::TypeInner::Struct { span, members };
let ty_handle = module.types.insert(
crate::Type {
@ -4111,7 +4101,7 @@ impl<I: Iterator<Item = u32>> Parser<I> {
self.span_from_with_op(start),
);
if block_decor == Some(Block { buffer: true }) {
if is_storage_buffer {
self.lookup_storage_buffer_types
.insert(ty_handle, storage_access);
}

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

@ -629,7 +629,6 @@ mod type_inner_tests {
crate::Type {
name: Some("MyType1".to_string()),
inner: crate::TypeInner::Struct {
top_level: true,
members: vec![],
span: 0,
},
@ -640,7 +639,6 @@ mod type_inner_tests {
crate::Type {
name: Some("MyType2".to_string()),
inner: crate::TypeInner::Struct {
top_level: true,
members: vec![],
span: 0,
},
@ -3981,7 +3979,6 @@ impl Parser {
let mut binding = None;
// Perspective is the default qualifier.
let mut stage = None;
let mut is_block = false;
let mut workgroup_size = [0u32; 3];
let mut early_depth_test = None;
@ -3995,9 +3992,6 @@ impl Parser {
bind_index = Some(parse_non_negative_sint_literal(lexer, 4)?);
lexer.expect(Token::Paren(')'))?;
}
("block", _) => {
is_block = true;
}
("group", _) => {
lexer.expect(Token::Paren('('))?;
bind_group = Some(parse_non_negative_sint_literal(lexer, 4)?);
@ -4077,11 +4071,7 @@ impl Parser {
let ty = module.types.insert(
crate::Type {
name: Some(name.to_string()),
inner: crate::TypeInner::Struct {
top_level: is_block,
members,
span,
},
inner: crate::TypeInner::Struct { members, span },
},
type_span,
);

3
src/front/wgsl/tests.rs
View File

@ -157,7 +157,7 @@ fn parse_type_cast() {
fn parse_struct() {
parse_str(
"
[[block]] struct Foo { x: i32; };
struct Foo { x: i32; };
struct Bar {
[[size(16)]] x: vec2<i32>;
[[align(16)]] y: f32;
@ -397,7 +397,6 @@ fn parse_struct_instantiation() {
fn parse_array_length() {
parse_str(
"
[[block]]
struct Foo {
data: [[stride(4)]] array<u32>;
}; // this is used as both input and output for convenience

14
src/lib.rs
View File

@ -627,11 +627,10 @@ pub enum TypeInner {
/// An `Array` is [`SIZED`] unless its `size` is [`Dynamic`].
/// Dynamically-sized arrays may only appear in a few situations:
///
/// - They may appear as the last member of a [`Struct`] whose `top_level`
/// flag is set.
/// - They may appear as the last member of a [`Struct`].
///
/// - They may appear as the base type of a [`Pointer`]. An
/// [`AccessIndex`] expression referring to a top-level struct's final
/// [`AccessIndex`] expression referring to a struct's final
/// unsized array member would have such a pointer type. However, such
/// pointer types may only appear as the types of such intermediate
/// expressions. They are not [`DATA`], and cannot be stored in
@ -655,20 +654,13 @@ pub enum TypeInner {
/// `DATA` as well.
///
/// Member types must be [`SIZED`], except for the final member of a
/// top-level struct, which may be a dynamically sized [`Array`]. The
/// struct, which may be a dynamically sized [`Array`]. The
/// `Struct` type itself is `SIZED` when all its members are `SIZED`.
///
/// When `top_level` is true, this `Struct` represents the contents of a
/// buffer resource occupying a single binding slot in a shader's resource
/// interface. Top-level `Struct`s may not be used as members of any other
/// struct, or as array elements.
///
/// [`DATA`]: crate::valid::TypeFlags::DATA
/// [`SIZED`]: crate::valid::TypeFlags::SIZED
/// [`Array`]: TypeInner::Array
Struct {
/// This struct serves as the type of a binding slot in a shader's resource interface.
top_level: bool,
members: Vec<StructMember>,
//TODO: should this be unaligned?
span: u32,

6
src/proc/layouter.rs
View File

@ -118,11 +118,7 @@ impl Layouter {
return Err(InvalidBaseType(base));
},
},
Ti::Struct {
top_level: _,
span,
ref members,
} => {
Ti::Struct { span, ref members } => {
let mut alignment = Alignment::new(1).unwrap();
for member in members {
alignment = if member.ty < ty_handle {

17
src/valid/interface.rs
View File

@ -298,11 +298,7 @@ impl VaryingContext<'_> {
None => {
match self.types[self.ty].inner {
//TODO: check the member types
crate::TypeInner::Struct {
top_level: false,
ref members,
..
} => {
crate::TypeInner::Struct { ref members, .. } => {
for (index, member) in members.iter().enumerate() {
self.ty = member.ty;
let span_context = self.types.get_span_context(self.ty);
@ -346,10 +342,7 @@ impl super::Validator {
return Err(GlobalVariableError::Alignment(ty_handle, disalignment));
}
}
(
TypeFlags::DATA | TypeFlags::HOST_SHARED | TypeFlags::TOP_LEVEL,
true,
)
(TypeFlags::DATA | TypeFlags::HOST_SHARED, true)
}
crate::StorageClass::Uniform => {
if let Err((ty_handle, disalignment)) = type_info.uniform_layout {
@ -358,11 +351,7 @@ impl super::Validator {
}
}
(
TypeFlags::DATA
| TypeFlags::COPY
| TypeFlags::SIZED
| TypeFlags::HOST_SHARED
| TypeFlags::TOP_LEVEL,
TypeFlags::DATA | TypeFlags::COPY | TypeFlags::SIZED | TypeFlags::HOST_SHARED,
true,
)
}

20
src/valid/type.rs
View File

@ -47,9 +47,6 @@ bitflags::bitflags! {
/// Can be used for host-shareable structures.
const HOST_SHARED = 0x10;
/// This is a top-level host-shareable type.
const TOP_LEVEL = 0x20;
/// This type can be passed as a function argument.
const ARGUMENT = 0x40;
}
@ -113,8 +110,6 @@ pub enum TypeError {
size: u32,
span: u32,
},
#[error("The composite type contains a top-level structure")]
NestedTopLevel,
}
// Only makes sense if `flags.contains(HOST_SHARED)`
@ -330,9 +325,6 @@ impl super::Validator {
if !base_info.flags.contains(TypeFlags::DATA | TypeFlags::SIZED) {
return Err(TypeError::InvalidArrayBaseType(base));
}
if base_info.flags.contains(TypeFlags::TOP_LEVEL) {
return Err(TypeError::NestedTopLevel);
}
let base_size = types[base].inner.span(constants);
if stride < base_size {
@ -443,11 +435,7 @@ impl super::Validator {
storage_layout,
}
}
Ti::Struct {
top_level,
ref members,
span,
} => {
Ti::Struct { ref members, span } => {
let mut ti = TypeInfo::new(
TypeFlags::DATA
| TypeFlags::SIZED
@ -467,9 +455,6 @@ impl super::Validator {
if !base_info.flags.contains(TypeFlags::DATA) {
return Err(TypeError::InvalidData(member.ty));
}
if base_info.flags.contains(TypeFlags::TOP_LEVEL) {
return Err(TypeError::NestedTopLevel);
}
if !base_info.flags.contains(TypeFlags::HOST_SHARED) {
if ti.uniform_layout.is_ok() {
ti.uniform_layout = Err((member.ty, Disalignment::NonHostShareable));
@ -535,9 +520,6 @@ impl super::Validator {
}
}
}
if top_level {
ti.flags |= TypeFlags::TOP_LEVEL;
}
let alignment = self.layouter[handle].alignment.get();
if span % alignment != 0 {

1
tests/in/access.wgsl
View File

@ -1,6 +1,5 @@
// This snapshot tests accessing various containers, dereferencing pointers.
[[block]]
struct Bar {
matrix: mat4x4<f32>;
atom: atomic<i32>;

2
tests/in/boids.wgsl
View File

@ -5,7 +5,6 @@ struct Particle {
vel : vec2<f32>;
};
[[block]]
struct SimParams {
deltaT : f32;
rule1Distance : f32;
@ -16,7 +15,6 @@ struct SimParams {
rule3Scale : f32;
};
[[block]]
struct Particles {
particles : [[stride(16)]] array<Particle>;
};

1
tests/in/bounds-check-restrict.wgsl
View File

@ -1,6 +1,5 @@
// Tests for `naga::back::BoundsCheckPolicy::Restrict`.
[[block]]
struct Globals {
a: array<f32, 10>;
v: vec4<f32>;

1
tests/in/bounds-check-zero.wgsl
View File

@ -1,6 +1,5 @@
// Tests for `naga::back::BoundsCheckPolicy::ReadZeroSkipWrite`.
[[block]]
struct Globals {
a: array<f32, 10>;
v: vec4<f32>;

1
tests/in/collatz.wgsl
View File

@ -1,4 +1,3 @@
[[block]]
struct PrimeIndices {
data: [[stride(4)]] array<u32>;
}; // this is used as both input and output for convenience

1
tests/in/extra.wgsl
View File

@ -1,4 +1,3 @@
[[block]]
struct PushConstants {
index: u32;
double: vec2<f64>;

1
tests/in/pointers.wgsl
View File

@ -4,7 +4,6 @@ fn f() {
*px = 10;
}
[[block]]
struct DynamicArray {
arr: array<u32>;
};

2
tests/in/policy-mix.wgsl
View File

@ -2,13 +2,11 @@
// implemented separately.
// Storage and Uniform storage classes
[[block]]
struct InStorage {
a: array<vec4<f32>, 10>;
};
[[group(0), binding(0)]] var<storage> in_storage: InStorage;
[[block]]
struct InUniform {
a: array<vec4<f32>, 20>;
};

2
tests/in/shadow.wgsl
View File

@ -1,4 +1,3 @@
[[block]]
struct Globals {
num_lights: vec4<u32>;
};
@ -12,7 +11,6 @@ struct Light {
color: vec4<f32>;
};
[[block]]
struct Lights {
data: [[stride(96)]] array<Light>;
};

1
tests/in/skybox.wgsl
View File

@ -3,7 +3,6 @@ struct VertexOutput {
[[location(0)]] uv: vec3<f32>;
};
[[block]]
struct Data {
proj_inv: mat4x4<f32>;
view: mat4x4<f32>;

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

@ -5,7 +5,7 @@ precision highp int;
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(std430) buffer Bar_block_0Cs {
layout(std430) buffer Bar_block_0Compute {
mat4x4 matrix;
int atom;
uvec2 arr[2];

2
tests/out/glsl/access.foo.Vertex.glsl
View File

@ -3,7 +3,7 @@
precision highp float;
precision highp int;
layout(std430) buffer Bar_block_0Vs {
layout(std430) buffer Bar_block_0Vertex {
mat4x4 matrix;
int atom;
uvec2 arr[2];

10
tests/out/glsl/boids.main.Compute.glsl
View File

@ -9,8 +9,7 @@ struct Particle {
vec2 pos;
vec2 vel;
};
uniform SimParams_block_0Cs {
struct SimParams {
float deltaT;
float rule1Distance;
float rule2Distance;
@ -18,13 +17,14 @@ uniform SimParams_block_0Cs {
float rule1Scale;
float rule2Scale;
float rule3Scale;
} _group_0_binding_0;
};
uniform SimParams_block_0Compute { SimParams _group_0_binding_0; };
layout(std430) readonly buffer Particles_block_1Cs {
layout(std430) readonly buffer Particles_block_1Compute {
Particle particles[];
} _group_0_binding_1;
layout(std430) buffer Particles_block_2Cs {
layout(std430) buffer Particles_block_2Compute {
Particle particles[];
} _group_0_binding_2;

1
tests/out/glsl/interpolate.frag_main.Fragment.glsl
View File

@ -9,7 +9,6 @@ struct FragmentInput {
float perspective_centroid;
float perspective_sample;
};
flat in uint _vs2fs_location0;
noperspective in float _vs2fs_location1;
noperspective centroid in vec2 _vs2fs_location2;

25
tests/out/glsl/interpolate.main.Fragment.glsl
View File

@ -1,25 +0,0 @@
#version 400 core
struct FragmentInput {
vec4 position;
uint flat_;
float linear;
vec2 linear_centroid;
vec3 linear_sample;
vec4 perspective;
float perspective_centroid;
float perspective_sample;
};
flat in uint _vs2fs_location0;
noperspective in float _vs2fs_location1;
noperspective centroid in vec2 _vs2fs_location2;
noperspective sample in vec3 _vs2fs_location3;
smooth in vec4 _vs2fs_location4;
smooth centroid in float _vs2fs_location5;
smooth sample in float _vs2fs_location6;
void main() {
FragmentInput val = FragmentInput(gl_FragCoord, _vs2fs_location0, _vs2fs_location1, _vs2fs_location2, _vs2fs_location3, _vs2fs_location4, _vs2fs_location5, _vs2fs_location6);
return;
}

42
tests/out/glsl/interpolate.main.Vertex.glsl
View File

@ -1,42 +0,0 @@
#version 400 core
struct FragmentInput {
vec4 position;
uint flat_;
float linear;
vec2 linear_centroid;
vec3 linear_sample;
vec4 perspective;
float perspective_centroid;
float perspective_sample;
};
flat out uint _vs2fs_location0;
noperspective out float _vs2fs_location1;
noperspective centroid out vec2 _vs2fs_location2;
noperspective sample out vec3 _vs2fs_location3;
smooth out vec4 _vs2fs_location4;
smooth centroid out float _vs2fs_location5;
smooth sample out float _vs2fs_location6;
void main() {
FragmentInput out_;
out_.position = vec4(2.0, 4.0, 5.0, 6.0);
out_.flat_ = 8u;
out_.linear = 27.0;
out_.linear_centroid = vec2(64.0, 125.0);
out_.linear_sample = vec3(216.0, 343.0, 512.0);
out_.perspective = vec4(729.0, 1000.0, 1331.0, 1728.0);
out_.perspective_centroid = 2197.0;
out_.perspective_sample = 2744.0;
FragmentInput _e30 = out_;
gl_Position = _e30.position;
_vs2fs_location0 = _e30.flat_;
_vs2fs_location1 = _e30.linear;
_vs2fs_location2 = _e30.linear_centroid;
_vs2fs_location3 = _e30.linear_sample;
_vs2fs_location4 = _e30.perspective;
_vs2fs_location5 = _e30.perspective_centroid;
_vs2fs_location6 = _e30.perspective_sample;
return;
}

1
tests/out/glsl/interpolate.vert_main.Vertex.glsl
View File

@ -9,7 +9,6 @@ struct FragmentInput {
float perspective_centroid;
float perspective_sample;
};
flat out uint _vs2fs_location0;
noperspective out float _vs2fs_location1;
noperspective centroid out vec2 _vs2fs_location2;

1
tests/out/glsl/operators.main.Compute.glsl
View File

@ -10,7 +10,6 @@ struct Foo {
int b;
};
vec4 builtins() {
int s1_ = (true ? 1 : 0);
vec4 s2_ = (true ? vec4(1.0, 1.0, 1.0, 1.0) : vec4(0.0, 0.0, 0.0, 0.0));

22
tests/out/glsl/quad-vert.main.Vertex.glsl
View File

@ -3,21 +3,21 @@
precision highp float;
precision highp int;
struct type_9 {
vec2 member;
vec4 gen_gl_Position;
};
vec2 v_uv = vec2(0.0, 0.0);
vec2 a_uv_1 = vec2(0.0, 0.0);
struct gen_gl_PerVertex_block_0Vs {
struct gen_gl_PerVertex {
vec4 gen_gl_Position;
float gen_gl_PointSize;
float gen_gl_ClipDistance[1];
float gen_gl_CullDistance[1];
} perVertexStruct;
};
struct type_9 {
vec2 member;
vec4 gen_gl_Position;
};
vec2 v_uv = vec2(0.0, 0.0);
vec2 a_uv_1 = vec2(0.0, 0.0);
gen_gl_PerVertex perVertexStruct;
vec2 a_pos_1 = vec2(0.0, 0.0);

1
tests/out/glsl/quad.frag_main.Fragment.glsl
View File

@ -7,7 +7,6 @@ struct VertexOutput {
vec2 uv;
vec4 position;
};
uniform highp sampler2D _group_0_binding_0;
smooth in vec2 _vs2fs_location0;

1
tests/out/glsl/quad.fs_extra.Fragment.glsl
View File

@ -7,7 +7,6 @@ struct VertexOutput {
vec2 uv;
vec4 position;
};
layout(location = 0) out vec4 _fs2p_location0;
void main() {

1
tests/out/glsl/quad.vert_main.Vertex.glsl
View File

@ -7,7 +7,6 @@ struct VertexOutput {
vec2 uv;
vec4 position;
};
layout(location = 0) in vec2 _p2vs_location0;
layout(location = 1) in vec2 _p2vs_location1;
smooth out vec2 _vs2fs_location0;

10
tests/out/glsl/shadow.fs_main.Fragment.glsl
View File

@ -3,17 +3,17 @@
precision highp float;
precision highp int;
struct Globals {
uvec4 num_lights;
};
struct Light {
mat4x4 proj;
vec4 pos;
vec4 color;
};
uniform Globals_block_0Fragment { Globals _group_0_binding_0; };
uniform Globals_block_0Fs {
uvec4 num_lights;
} _group_0_binding_0;
layout(std430) readonly buffer Lights_block_1Fs {
layout(std430) readonly buffer Lights_block_1Fragment {
Light data[];
} _group_0_binding_1;

5
tests/out/glsl/skybox.fs_main.Fragment.glsl
View File

@ -7,7 +7,10 @@ struct VertexOutput {
vec4 position;
vec3 uv;
};
struct Data {
mat4x4 proj_inv;
mat4x4 view;
};
layout(binding = 0) uniform highp samplerCube _group_0_binding_1;
layout(location = 0) smooth in vec3 _vs2fs_location0;

6
tests/out/glsl/skybox.vs_main.Vertex.glsl
View File

@ -7,11 +7,11 @@ struct VertexOutput {
vec4 position;
vec3 uv;
};
layout(std140, binding = 0) uniform Data_block_0Vs {
struct Data {
mat4x4 proj_inv;
mat4x4 view;
} _group_0_binding_0;
};
layout(std140, binding = 0) uniform Data_block_0Vertex { Data _group_0_binding_0; };
layout(location = 0) smooth out vec3 _vs2fs_location0;

1
tests/out/ir/collatz.ron
View File

@ -18,7 +18,6 @@
(
name: Some("PrimeIndices"),
inner: Struct(
top_level: true,
members: [
(
name: Some("data"),

3
tests/out/ir/shadow.ron
View File

@ -94,7 +94,6 @@
(
name: Some("Globals"),
inner: Struct(
top_level: true,
members: [
(
name: Some("num_lights"),
@ -138,7 +137,6 @@
(
name: Some("Light"),
inner: Struct(
top_level: false,
members: [
(
name: Some("proj"),
@ -173,7 +171,6 @@
(
name: Some("Lights"),
inner: Struct(
top_level: true,
members: [
(
name: Some("data"),

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

@ -26,7 +26,6 @@ OpName %90 "tmp"
OpName %92 "atomics"
OpDecorate %24 ArrayStride 8
OpDecorate %25 ArrayStride 8
OpDecorate %26 Block
OpMemberDecorate %26 0 Offset 0
OpMemberDecorate %26 0 ColMajor
OpMemberDecorate %26 0 MatrixStride 16
@ -36,6 +35,7 @@ OpMemberDecorate %26 3 Offset 88
OpDecorate %29 ArrayStride 4
OpDecorate %30 DescriptorSet 0
OpDecorate %30 Binding 0
OpDecorate %26 Block
OpDecorate %42 BuiltIn VertexIndex
OpDecorate %45 BuiltIn Position
%2 = OpTypeVoid

545
tests/out/spv/boids.spvasm
View File

@ -1,13 +1,13 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 203
; Bound: 206
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %43 "main" %40
OpExecutionMode %43 LocalSize 64 1 1
OpEntryPoint GLCompute %44 "main" %41
OpExecutionMode %44 LocalSize 64 1 1
OpSource GLSL 450
OpName %3 "NUM_PARTICLES"
OpMemberName %16 0 "pos"
@ -24,23 +24,22 @@ OpName %17 "SimParams"
OpMemberName %19 0 "particles"
OpName %19 "Particles"
OpName %21 "params"
OpName %23 "particlesSrc"
OpName %25 "particlesDst"
OpName %26 "vPos"
OpName %28 "vVel"
OpName %29 "cMass"
OpName %30 "cVel"
OpName %31 "colVel"
OpName %32 "cMassCount"
OpName %34 "cVelCount"
OpName %35 "pos"
OpName %36 "vel"
OpName %37 "i"
OpName %40 "global_invocation_id"
OpName %43 "main"
OpName %24 "particlesSrc"
OpName %26 "particlesDst"
OpName %27 "vPos"
OpName %29 "vVel"
OpName %30 "cMass"
OpName %31 "cVel"
OpName %32 "colVel"
OpName %33 "cMassCount"
OpName %35 "cVelCount"
OpName %36 "pos"
OpName %37 "vel"
OpName %38 "i"
OpName %41 "global_invocation_id"
OpName %44 "main"
OpMemberDecorate %16 0 Offset 0
OpMemberDecorate %16 1 Offset 8
OpDecorate %17 Block
OpMemberDecorate %17 0 Offset 0
OpMemberDecorate %17 1 Offset 4
OpMemberDecorate %17 2 Offset 8
@ -49,16 +48,19 @@ OpMemberDecorate %17 4 Offset 16
OpMemberDecorate %17 5 Offset 20
OpMemberDecorate %17 6 Offset 24
OpDecorate %18 ArrayStride 16
OpDecorate %19 Block
OpMemberDecorate %19 0 Offset 0
OpDecorate %21 DescriptorSet 0
OpDecorate %21 Binding 0
OpDecorate %23 NonWritable
OpDecorate %23 DescriptorSet 0
OpDecorate %23 Binding 1
OpDecorate %25 DescriptorSet 0
OpDecorate %25 Binding 2
OpDecorate %40 BuiltIn GlobalInvocationId
OpDecorate %22 Block
OpMemberDecorate %22 0 Offset 0
OpDecorate %24 NonWritable
OpDecorate %24 DescriptorSet 0
OpDecorate %24 Binding 1
OpDecorate %19 Block
OpDecorate %26 DescriptorSet 0
OpDecorate %26 Binding 2
OpDecorate %19 Block
OpDecorate %41 BuiltIn GlobalInvocationId
%2 = OpTypeVoid
%4 = OpTypeInt 32 0
%3 = OpConstant %4 1500
@ -78,253 +80,256 @@ OpDecorate %40 BuiltIn GlobalInvocationId
%18 = OpTypeRuntimeArray %16
%19 = OpTypeStruct %18
%20 = OpTypeVector %4 3
%22 = OpTypePointer Uniform %17
%21 = OpVariable %22 Uniform
%24 = OpTypePointer StorageBuffer %19
%23 = OpVariable %24 StorageBuffer
%25 = OpVariable %24 StorageBuffer
%27 = OpTypePointer Function %15
%33 = OpTypePointer Function %8
%38 = OpTypePointer Function %4
%41 = OpTypePointer Input %20
%40 = OpVariable %41 Input
%44 = OpTypeFunction %2
%47 = OpTypeBool
%51 = OpTypePointer StorageBuffer %18
%52 = OpTypePointer StorageBuffer %16
%53 = OpTypePointer StorageBuffer %15
%82 = OpTypePointer Uniform %6
%96 = OpConstant %4 2
%110 = OpConstant %4 3
%145 = OpConstant %4 4
%151 = OpConstant %4 5
%157 = OpConstant %4 6
%174 = OpTypePointer Function %6
%43 = OpFunction %2 None %44
%39 = OpLabel
%37 = OpVariable %38 Function %9
%34 = OpVariable %33 Function %7
%30 = OpVariable %27 Function
%26 = OpVariable %27 Function
%35 = OpVariable %27 Function
%31 = OpVariable %27 Function
%28 = OpVariable %27 Function
%36 = OpVariable %27 Function
%32 = OpVariable %33 Function %7
%29 = OpVariable %27 Function
%42 = OpLoad %20 %40
OpBranch %45
%45 = OpLabel
%46 = OpCompositeExtract %4 %42 0
%48 = OpUGreaterThanEqual %47 %46 %3
OpSelectionMerge %49 None
OpBranchConditional %48 %50 %49
%50 = OpLabel
%22 = OpTypeStruct %17
%23 = OpTypePointer Uniform %22
%21 = OpVariable %23 Uniform
%25 = OpTypePointer StorageBuffer %19
%24 = OpVariable %25 StorageBuffer
%26 = OpVariable %25 StorageBuffer
%28 = OpTypePointer Function %15
%34 = OpTypePointer Function %8
%39 = OpTypePointer Function %4
%42 = OpTypePointer Input %20
%41 = OpVariable %42 Input
%45 = OpTypeFunction %2
%46 = OpTypePointer Uniform %17
%50 = OpTypeBool
%54 = OpTypePointer StorageBuffer %18
%55 = OpTypePointer StorageBuffer %16
%56 = OpTypePointer StorageBuffer %15
%85 = OpTypePointer Uniform %6
%99 = OpConstant %4 2
%113 = OpConstant %4 3
%148 = OpConstant %4 4
%154 = OpConstant %4 5
%160 = OpConstant %4 6
%177 = OpTypePointer Function %6
%44 = OpFunction %2 None %45
%40 = OpLabel
%38 = OpVariable %39 Function %9
%35 = OpVariable %34 Function %7
%31 = OpVariable %28 Function
%27 = OpVariable %28 Function
%36 = OpVariable %28 Function
%32 = OpVariable %28 Function
%29 = OpVariable %28 Function
%37 = OpVariable %28 Function
%33 = OpVariable %34 Function %7
%30 = OpVariable %28 Function
%43 = OpLoad %20 %41
%47 = OpAccessChain %46 %21 %9
OpBranch %48
%48 = OpLabel
%49 = OpCompositeExtract %4 %43 0
%51 = OpUGreaterThanEqual %50 %49 %3
OpSelectionMerge %52 None
OpBranchConditional %51 %53 %52
%53 = OpLabel
OpReturn
%49 = OpLabel
%54 = OpAccessChain %53 %23 %9 %46 %9
%55 = OpLoad %15 %54
OpStore %26 %55
%56 = OpAccessChain %53 %23 %9 %46 %11
%57 = OpLoad %15 %56
OpStore %28 %57
%58 = OpCompositeConstruct %15 %5 %5
OpStore %29 %58
%59 = OpCompositeConstruct %15 %5 %5
OpStore %30 %59
%60 = OpCompositeConstruct %15 %5 %5
OpStore %31 %60
OpBranch %61
%61 = OpLabel
OpLoopMerge %62 %64 None
OpBranch %63
%63 = OpLabel
%65 = OpLoad %4 %37
%66 = OpUGreaterThanEqual %47 %65 %3
OpSelectionMerge %67 None
OpBranchConditional %66 %68 %67
%68 = OpLabel
OpBranch %62
%67 = OpLabel
%69 = OpLoad %4 %37
%70 = OpIEqual %47 %69 %46
OpSelectionMerge %71 None
OpBranchConditional %70 %72 %71
%72 = OpLabel
OpBranch %64
%71 = OpLabel
%73 = OpLoad %4 %37
%74 = OpAccessChain %53 %23 %9 %73 %9
%75 = OpLoad %15 %74
OpStore %35 %75
%76 = OpLoad %4 %37
%77 = OpAccessChain %53 %23 %9 %76 %11
%78 = OpLoad %15 %77
OpStore %36 %78
%79 = OpLoad %15 %35
%80 = OpLoad %15 %26
%81 = OpExtInst %6 %1 Distance %79 %80
%83 = OpAccessChain %82 %21 %11
%84 = OpLoad %6 %83
%85 = OpFOrdLessThan %47 %81 %84
OpSelectionMerge %86 None
OpBranchConditional %85 %87 %86
%87 = OpLabel
%88 = OpLoad %15 %29
%89 = OpLoad %15 %35
%90 = OpFAdd %15 %88 %89
OpStore %29 %90
%91 = OpLoad %8 %32
%92 = OpIAdd %8 %91 %10
OpStore %32 %92
OpBranch %86
%86 = OpLabel
%93 = OpLoad %15 %35
%94 = OpLoad %15 %26
%95 = OpExtInst %6 %1 Distance %93 %94
%97 = OpAccessChain %82 %21 %96
%98 = OpLoad %6 %97
%99 = OpFOrdLessThan %47 %95 %98
OpSelectionMerge %100 None
OpBranchConditional %99 %101 %100
%101 = OpLabel
%102 = OpLoad %15 %31
%103 = OpLoad %15 %35
%104 = OpLoad %15 %26
%105 = OpFSub %15 %103 %104
%106 = OpFSub %15 %102 %105
OpStore %31 %106
OpBranch %100
%100 = OpLabel
%107 = OpLoad %15 %35
%108 = OpLoad %15 %26
%109 = OpExtInst %6 %1 Distance %107 %108
%111 = OpAccessChain %82 %21 %110
%112 = OpLoad %6 %111
%113 = OpFOrdLessThan %47 %109 %112
OpSelectionMerge %114 None
OpBranchConditional %113 %115 %114
%115 = OpLabel
%116 = OpLoad %15 %30
%117 = OpLoad %15 %36
%118 = OpFAdd %15 %116 %117
OpStore %30 %118
%119 = OpLoad %8 %34
%120 = OpIAdd %8 %119 %10
OpStore %34 %120
OpBranch %114
%114 = OpLabel
%52 = OpLabel
%57 = OpAccessChain %56 %24 %9 %49 %9
%58 = OpLoad %15 %57
OpStore %27 %58
%59 = OpAccessChain %56 %24 %9 %49 %11
%60 = OpLoad %15 %59
OpStore %29 %60
%61 = OpCompositeConstruct %15 %5 %5
OpStore %30 %61
%62 = OpCompositeConstruct %15 %5 %5
OpStore %31 %62
%63 = OpCompositeConstruct %15 %5 %5
OpStore %32 %63
OpBranch %64
%64 = OpLabel
%121 = OpLoad %4 %37
%122 = OpIAdd %4 %121 %11
OpStore %37 %122
OpBranch %61
%62 = OpLabel
%123 = OpLoad %8 %32
%124 = OpSGreaterThan %47 %123 %7
OpSelectionMerge %125 None
OpBranchConditional %124 %126 %125
%126 = OpLabel
%127 = OpLoad %15 %29
%128 = OpLoad %8 %32
%129 = OpConvertSToF %6 %128
%130 = OpCompositeConstruct %15 %129 %129
%131 = OpFDiv %15 %127 %130
%132 = OpLoad %15 %26
%133 = OpFSub %15 %131 %132
OpStore %29 %133
OpBranch %125
%125 = OpLabel
%134 = OpLoad %8 %34
%135 = OpSGreaterThan %47 %134 %7
OpSelectionMerge %136 None
OpBranchConditional %135 %137 %136
%137 = OpLabel
%138 = OpLoad %15 %30
%139 = OpLoad %8 %34
%140 = OpConvertSToF %6 %139
%141 = OpCompositeConstruct %15 %140 %140
%142 = OpFDiv %15 %138 %141
OpStore %30 %142
OpBranch %136
%136 = OpLabel
%143 = OpLoad %15 %28
%144 = OpLoad %15 %29
%146 = OpAccessChain %82 %21 %145
%147 = OpLoad %6 %146
%148 = OpVectorTimesScalar %15 %144 %147
%149 = OpFAdd %15 %143 %148
%150 = OpLoad %15 %31
%152 = OpAccessChain %82 %21 %151
%153 = OpLoad %6 %152
%154 = OpVectorTimesScalar %15 %150 %153
%155 = OpFAdd %15 %149 %154
%156 = OpLoad %15 %30
%158 = OpAccessChain %82 %21 %157
%159 = OpLoad %6 %158
%160 = OpVectorTimesScalar %15 %156 %159
%161 = OpFAdd %15 %155 %160
OpStore %28 %161
%162 = OpLoad %15 %28
%163 = OpExtInst %15 %1 Normalize %162
%164 = OpLoad %15 %28
%165 = OpExtInst %6 %1 Length %164
%166 = OpExtInst %6 %1 FClamp %165 %5 %12
%167 = OpVectorTimesScalar %15 %163 %166
OpStore %28 %167
%168 = OpLoad %15 %26
%169 = OpLoad %15 %28
%170 = OpAccessChain %82 %21 %9
%171 = OpLoad %6 %170
%172 = OpVectorTimesScalar %15 %169 %171
%173 = OpFAdd %15 %168 %172
OpStore %26 %173
%175 = OpAccessChain %174 %26 %9
%176 = OpLoad %6 %175
%177 = OpFOrdLessThan %47 %176 %13
OpSelectionMerge %178 None
OpBranchConditional %177 %179 %178
%179 = OpLabel
%180 = OpAccessChain %174 %26 %9
OpStore %180 %14
OpBranch %178
%178 = OpLabel
%181 = OpAccessChain %174 %26 %9
%182 = OpLoad %6 %181
%183 = OpFOrdGreaterThan %47 %182 %14
OpSelectionMerge %184 None
OpBranchConditional %183 %185 %184
%185 = OpLabel
%186 = OpAccessChain %174 %26 %9
OpStore %186 %13
OpBranch %184
%184 = OpLabel
%187 = OpAccessChain %174 %26 %11
%188 = OpLoad %6 %187
%189 = OpFOrdLessThan %47 %188 %13
OpSelectionMerge %190 None
OpBranchConditional %189 %191 %190
%191 = OpLabel
%192 = OpAccessChain %174 %26 %11
OpStore %192 %14
OpBranch %190
%190 = OpLabel
%193 = OpAccessChain %174 %26 %11
%194 = OpLoad %6 %193
%195 = OpFOrdGreaterThan %47 %194 %14
OpSelectionMerge %196 None
OpBranchConditional %195 %197 %196
%197 = OpLabel
%198 = OpAccessChain %174 %26 %11
OpStore %198 %13
OpBranch %196
%196 = OpLabel
%199 = OpLoad %15 %26
%200 = OpAccessChain %53 %25 %9 %46 %9
OpStore %200 %199
%201 = OpLoad %15 %28
%202 = OpAccessChain %53 %25 %9 %46 %11
OpStore %202 %201
OpLoopMerge %65 %67 None
OpBranch %66
%66 = OpLabel
%68 = OpLoad %4 %38
%69 = OpUGreaterThanEqual %50 %68 %3
OpSelectionMerge %70 None
OpBranchConditional %69 %71 %70
%71 = OpLabel
OpBranch %65
%70 = OpLabel
%72 = OpLoad %4 %38
%73 = OpIEqual %50 %72 %49
OpSelectionMerge %74 None
OpBranchConditional %73 %75 %74
%75 = OpLabel
OpBranch %67
%74 = OpLabel
%76 = OpLoad %4 %38
%77 = OpAccessChain %56 %24 %9 %76 %9
%78 = OpLoad %15 %77
OpStore %36 %78
%79 = OpLoad %4 %38
%80 = OpAccessChain %56 %24 %9 %79 %11
%81 = OpLoad %15 %80
OpStore %37 %81
%82 = OpLoad %15 %36
%83 = OpLoad %15 %27
%84 = OpExtInst %6 %1 Distance %82 %83
%86 = OpAccessChain %85 %47 %11
%87 = OpLoad %6 %86
%88 = OpFOrdLessThan %50 %84 %87
OpSelectionMerge %89 None
OpBranchConditional %88 %90 %89
%90 = OpLabel
%91 = OpLoad %15 %30
%92 = OpLoad %15 %36
%93 = OpFAdd %15 %91 %92
OpStore %30 %93
%94 = OpLoad %8 %33
%95 = OpIAdd %8 %94 %10
OpStore %33 %95
OpBranch %89
%89 = OpLabel
%96 = OpLoad %15 %36
%97 = OpLoad %15 %27
%98 = OpExtInst %6 %1 Distance %96 %97
%100 = OpAccessChain %85 %47 %99
%101 = OpLoad %6 %100
%102 = OpFOrdLessThan %50 %98 %101
OpSelectionMerge %103 None
OpBranchConditional %102 %104 %103
%104 = OpLabel
%105 = OpLoad %15 %32
%106 = OpLoad %15 %36
%107 = OpLoad %15 %27
%108 = OpFSub %15 %106 %107
%109 = OpFSub %15 %105 %108
OpStore %32 %109
OpBranch %103
%103 = OpLabel
%110 = OpLoad %15 %36
%111 = OpLoad %15 %27
%112 = OpExtInst %6 %1 Distance %110 %111
%114 = OpAccessChain %85 %47 %113
%115 = OpLoad %6 %114
%116 = OpFOrdLessThan %50 %112 %115
OpSelectionMerge %117 None
OpBranchConditional %116 %118 %117
%118 = OpLabel
%119 = OpLoad %15 %31
%120 = OpLoad %15 %37
%121 = OpFAdd %15 %119 %120
OpStore %31 %121
%122 = OpLoad %8 %35
%123 = OpIAdd %8 %122 %10
OpStore %35 %123
OpBranch %117
%117 = OpLabel
OpBranch %67
%67 = OpLabel
%124 = OpLoad %4 %38
%125 = OpIAdd %4 %124 %11
OpStore %38 %125
OpBranch %64
%65 = OpLabel
%126 = OpLoad %8 %33
%127 = OpSGreaterThan %50 %126 %7
OpSelectionMerge %128 None
OpBranchConditional %127 %129 %128
%129 = OpLabel
%130 = OpLoad %15 %30
%131 = OpLoad %8 %33
%132 = OpConvertSToF %6 %131
%133 = OpCompositeConstruct %15 %132 %132
%134 = OpFDiv %15 %130 %133
%135 = OpLoad %15 %27
%136 = OpFSub %15 %134 %135
OpStore %30 %136
OpBranch %128
%128 = OpLabel
%137 = OpLoad %8 %35
%138 = OpSGreaterThan %50 %137 %7
OpSelectionMerge %139 None
OpBranchConditional %138 %140 %139
%140 = OpLabel
%141 = OpLoad %15 %31
%142 = OpLoad %8 %35
%143 = OpConvertSToF %6 %142
%144 = OpCompositeConstruct %15 %143 %143
%145 = OpFDiv %15 %141 %144
OpStore %31 %145
OpBranch %139
%139 = OpLabel
%146 = OpLoad %15 %29
%147 = OpLoad %15 %30
%149 = OpAccessChain %85 %47 %148
%150 = OpLoad %6 %149
%151 = OpVectorTimesScalar %15 %147 %150
%152 = OpFAdd %15 %146 %151
%153 = OpLoad %15 %32
%155 = OpAccessChain %85 %47 %154
%156 = OpLoad %6 %155
%157 = OpVectorTimesScalar %15 %153 %156
%158 = OpFAdd %15 %152 %157
%159 = OpLoad %15 %31
%161 = OpAccessChain %85 %47 %160
%162 = OpLoad %6 %161
%163 = OpVectorTimesScalar %15 %159 %162
%164 = OpFAdd %15 %158 %163
OpStore %29 %164
%165 = OpLoad %15 %29
%166 = OpExtInst %15 %1 Normalize %165
%167 = OpLoad %15 %29
%168 = OpExtInst %6 %1 Length %167
%169 = OpExtInst %6 %1 FClamp %168 %5 %12
%170 = OpVectorTimesScalar %15 %166 %169
OpStore %29 %170
%171 = OpLoad %15 %27
%172 = OpLoad %15 %29
%173 = OpAccessChain %85 %47 %9
%174 = OpLoad %6 %173
%175 = OpVectorTimesScalar %15 %172 %174
%176 = OpFAdd %15 %171 %175
OpStore %27 %176
%178 = OpAccessChain %177 %27 %9
%179 = OpLoad %6 %178
%180 = OpFOrdLessThan %50 %179 %13
OpSelectionMerge %181 None
OpBranchConditional %180 %182 %181
%182 = OpLabel
%183 = OpAccessChain %177 %27 %9
OpStore %183 %14
OpBranch %181
%181 = OpLabel
%184 = OpAccessChain %177 %27 %9
%185 = OpLoad %6 %184
%186 = OpFOrdGreaterThan %50 %185 %14
OpSelectionMerge %187 None
OpBranchConditional %186 %188 %187
%188 = OpLabel
%189 = OpAccessChain %177 %27 %9
OpStore %189 %13
OpBranch %187
%187 = OpLabel
%190 = OpAccessChain %177 %27 %11
%191 = OpLoad %6 %190
%192 = OpFOrdLessThan %50 %191 %13
OpSelectionMerge %193 None
OpBranchConditional %192 %194 %193
%194 = OpLabel
%195 = OpAccessChain %177 %27 %11
OpStore %195 %14
OpBranch %193
%193 = OpLabel
%196 = OpAccessChain %177 %27 %11
%197 = OpLoad %6 %196
%198 = OpFOrdGreaterThan %50 %197 %14
OpSelectionMerge %199 None
OpBranchConditional %198 %200 %199
%200 = OpLabel
%201 = OpAccessChain %177 %27 %11
OpStore %201 %13
OpBranch %199
%199 = OpLabel
%202 = OpLoad %15 %27
%203 = OpAccessChain %56 %26 %9 %49 %9
OpStore %203 %202
%204 = OpLoad %15 %29
%205 = OpAccessChain %56 %26 %9 %49 %11
OpStore %205 %204
OpReturn
OpFunctionEnd

2
tests/out/spv/bounds-check-restrict.spvasm
View File

@ -9,7 +9,6 @@ OpExtension "SPV_KHR_storage_buffer_storage_class"
OpMemoryModel Logical GLSL450
OpDecorate %10 ArrayStride 4
OpDecorate %13 ArrayStride 4
OpDecorate %14 Block
OpMemberDecorate %14 0 Offset 0
OpMemberDecorate %14 1 Offset 48
OpMemberDecorate %14 2 Offset 64
@ -18,6 +17,7 @@ OpMemberDecorate %14 2 MatrixStride 16
OpMemberDecorate %14 3 Offset 112
OpDecorate %15 DescriptorSet 0
OpDecorate %15 Binding 0
OpDecorate %14 Block
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpConstant %4 10

2
tests/out/spv/bounds-check-zero.spvasm
View File

@ -9,7 +9,6 @@ OpExtension "SPV_KHR_storage_buffer_storage_class"
OpMemoryModel Logical GLSL450
OpDecorate %10 ArrayStride 4
OpDecorate %13 ArrayStride 4
OpDecorate %14 Block
OpMemberDecorate %14 0 Offset 0
OpMemberDecorate %14 1 Offset 48
OpMemberDecorate %14 2 Offset 64
@ -18,6 +17,7 @@ OpMemberDecorate %14 2 MatrixStride 16
OpMemberDecorate %14 3 Offset 112
OpDecorate %15 DescriptorSet 0
OpDecorate %15 Binding 0
OpDecorate %14 Block
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpConstant %4 10

2
tests/out/spv/collatz.spvasm
View File

@ -19,10 +19,10 @@ OpName %18 "collatz_iterations"
OpName %45 "global_id"
OpName %48 "main"
OpDecorate %8 ArrayStride 4
OpDecorate %9 Block
OpMemberDecorate %9 0 Offset 0
OpDecorate %11 DescriptorSet 0
OpDecorate %11 Binding 0
OpDecorate %9 Block
OpDecorate %45 BuiltIn GlobalInvocationId
%2 = OpTypeVoid
%4 = OpTypeInt 32 0

1
tests/out/spv/extra.spvasm
View File

@ -9,7 +9,6 @@ OpCapability Geometry
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %25 "main" %17 %20 %23
OpExecutionMode %25 OriginUpperLeft
OpDecorate %10 Block
OpMemberDecorate %10 0 Offset 0
OpMemberDecorate %10 1 Offset 16
OpMemberDecorate %12 0 Offset 0

2
tests/out/spv/pointers.spvasm
View File

@ -20,10 +20,10 @@ OpName %33 "i"
OpName %34 "v"
OpName %35 "index_dynamic_array"
OpDecorate %7 ArrayStride 4
OpDecorate %8 Block
OpMemberDecorate %8 0 Offset 0
OpDecorate %11 DescriptorSet 0
OpDecorate %11 Binding 0
OpDecorate %8 Block
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpConstant %4 10

192
tests/out/spv/policy-mix.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 93
; Bound: 99
OpCapability Shader
OpCapability ImageQuery
OpCapability Linkage
@ -14,20 +14,18 @@ OpName %15 "InStorage"
OpMemberName %17 0 "a"
OpName %17 "InUniform"
OpName %23 "in_storage"
OpName %25 "in_uniform"
OpName %27 "image_2d_array"
OpName %29 "in_workgroup"
OpName %31 "in_private"
OpName %33 "in_function"
OpName %36 "c"
OpName %37 "i"
OpName %38 "l"
OpName %39 "mock_function"
OpName %26 "in_uniform"
OpName %29 "image_2d_array"
OpName %31 "in_workgroup"
OpName %33 "in_private"
OpName %35 "in_function"
OpName %38 "c"
OpName %39 "i"
OpName %40 "l"
OpName %41 "mock_function"
OpDecorate %14 ArrayStride 16
OpDecorate %15 Block
OpMemberDecorate %15 0 Offset 0
OpDecorate %16 ArrayStride 16
OpDecorate %17 Block
OpMemberDecorate %17 0 Offset 0
OpDecorate %19 ArrayStride 4
OpDecorate %20 ArrayStride 4
@ -35,10 +33,14 @@ OpDecorate %22 ArrayStride 16
OpDecorate %23 NonWritable
OpDecorate %23 DescriptorSet 0
OpDecorate %23 Binding 0
OpDecorate %25 DescriptorSet 0
OpDecorate %25 Binding 1
OpDecorate %27 DescriptorSet 0
OpDecorate %27 Binding 2
OpDecorate %24 Block
OpMemberDecorate %24 0 Offset 0
OpDecorate %26 DescriptorSet 0
OpDecorate %26 Binding 1
OpDecorate %27 Block
OpMemberDecorate %27 0 Offset 0
OpDecorate %29 DescriptorSet 0
OpDecorate %29 Binding 2
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpConstant %4 10
@ -60,81 +62,87 @@ OpDecorate %27 Binding 2
%20 = OpTypeArray %10 %7
%21 = OpTypeVector %4 2
%22 = OpTypeArray %13 %8
%24 = OpTypePointer StorageBuffer %15
%23 = OpVariable %24 StorageBuffer
%26 = OpTypePointer Uniform %17
%25 = OpVariable %26 Uniform
%28 = OpTypePointer UniformConstant %18
%27 = OpVariable %28 UniformConstant
%30 = OpTypePointer Workgroup %19
%29 = OpVariable %30 Workgroup
%32 = OpTypePointer Private %20
%31 = OpVariable %32 Private
%34 = OpTypePointer Function %22
%40 = OpTypeFunction %13 %21 %4 %4
%46 = OpTypePointer StorageBuffer %14
%47 = OpTypePointer StorageBuffer %13
%49 = OpTypeInt 32 0
%48 = OpConstant %49 0
%52 = OpTypePointer Uniform %16
%53 = OpTypePointer Uniform %13
%57 = OpTypeVector %4 3
%59 = OpTypeBool
%60 = OpConstantNull %13
%66 = OpTypeVector %59 3
%73 = OpTypePointer Workgroup %10
%74 = OpConstant %49 29
%80 = OpTypePointer Private %10
%81 = OpConstant %49 39
%87 = OpTypePointer Function %13
%88 = OpConstant %49 1
%39 = OpFunction %13 None %40
%36 = OpFunctionParameter %21
%37 = OpFunctionParameter %4
%38 = OpFunctionParameter %4
%35 = OpLabel
%33 = OpVariable %34 Function
%41 = OpLoad %18 %27
OpBranch %42
%42 = OpLabel
%43 = OpCompositeConstruct %13 %9 %11 %11 %12
%44 = OpCompositeConstruct %13 %11 %9 %11 %12
%45 = OpCompositeConstruct %22 %43 %44
OpStore %33 %45
%50 = OpAccessChain %47 %23 %48 %37
%51 = OpLoad %13 %50
%54 = OpAccessChain %53 %25 %48 %37
%55 = OpLoad %13 %54
%56 = OpFAdd %13 %51 %55
%58 = OpCompositeConstruct %57 %36 %37
%61 = OpImageQueryLevels %4 %41
%62 = OpULessThan %59 %38 %61
OpSelectionMerge %63 None
OpBranchConditional %62 %64 %63
%64 = OpLabel
%65 = OpImageQuerySizeLod %57 %41 %38
%67 = OpULessThan %66 %58 %65
%68 = OpAll %59 %67
OpBranchConditional %68 %69 %63
%24 = OpTypeStruct %15
%25 = OpTypePointer StorageBuffer %24
%23 = OpVariable %25 StorageBuffer
%27 = OpTypeStruct %17
%28 = OpTypePointer Uniform %27
%26 = OpVariable %28 Uniform
%30 = OpTypePointer UniformConstant %18
%29 = OpVariable %30 UniformConstant
%32 = OpTypePointer Workgroup %19
%31 = OpVariable %32 Workgroup
%34 = OpTypePointer Private %20
%33 = OpVariable %34 Private
%36 = OpTypePointer Function %22
%42 = OpTypeFunction %13 %21 %4 %4
%43 = OpTypePointer StorageBuffer %15
%45 = OpTypeInt 32 0
%44 = OpConstant %45 0
%47 = OpTypePointer Uniform %17
%54 = OpTypePointer StorageBuffer %14
%55 = OpTypePointer StorageBuffer %13
%58 = OpTypePointer Uniform %16
%59 = OpTypePointer Uniform %13
%63 = OpTypeVector %4 3
%65 = OpTypeBool
%66 = OpConstantNull %13
%72 = OpTypeVector %65 3
%79 = OpTypePointer Workgroup %10
%80 = OpConstant %45 29
%86 = OpTypePointer Private %10
%87 = OpConstant %45 39
%93 = OpTypePointer Function %13
%94 = OpConstant %45 1
%41 = OpFunction %13 None %42
%38 = OpFunctionParameter %21
%39 = OpFunctionParameter %4
%40 = OpFunctionParameter %4
%37 = OpLabel
%35 = OpVariable %36 Function
%46 = OpAccessChain %43 %23 %44
%48 = OpAccessChain %47 %26 %44
%49 = OpLoad %18 %29
OpBranch %50
%50 = OpLabel
%51 = OpCompositeConstruct %13 %9 %11 %11 %12
%52 = OpCompositeConstruct %13 %11 %9 %11 %12
%53 = OpCompositeConstruct %22 %51 %52
OpStore %35 %53
%56 = OpAccessChain %55 %46 %44 %39
%57 = OpLoad %13 %56
%60 = OpAccessChain %59 %48 %44 %39
%61 = OpLoad %13 %60
%62 = OpFAdd %13 %57 %61
%64 = OpCompositeConstruct %63 %38 %39
%67 = OpImageQueryLevels %4 %49
%68 = OpULessThan %65 %40 %67
OpSelectionMerge %69 None
OpBranchConditional %68 %70 %69
%70 = OpLabel
%71 = OpImageQuerySizeLod %63 %49 %40
%73 = OpULessThan %72 %64 %71
%74 = OpAll %65 %73
OpBranchConditional %74 %75 %69
%75 = OpLabel
%76 = OpImageFetch %13 %49 %64 Lod %40
OpBranch %69
%69 = OpLabel
%70 = OpImageFetch %13 %41 %58 Lod %38
OpBranch %63
%63 = OpLabel
%71 = OpPhi %13 %60 %42 %60 %64 %70 %69
%72 = OpFAdd %13 %56 %71
%75 = OpExtInst %49 %1 UMin %37 %74
%76 = OpAccessChain %73 %29 %75
%77 = OpLoad %10 %76
%78 = OpCompositeConstruct %13 %77 %77 %77 %77
%79 = OpFAdd %13 %72 %78
%82 = OpExtInst %49 %1 UMin %37 %81
%83 = OpAccessChain %80 %31 %82
%84 = OpLoad %10 %83
%85 = OpCompositeConstruct %13 %84 %84 %84 %84
%86 = OpFAdd %13 %79 %85
%89 = OpExtInst %49 %1 UMin %37 %88
%90 = OpAccessChain %87 %33 %89
%91 = OpLoad %13 %90
%92 = OpFAdd %13 %86 %91
OpReturnValue %92
%77 = OpPhi %13 %66 %50 %66 %70 %76 %75
%78 = OpFAdd %13 %62 %77
%81 = OpExtInst %45 %1 UMin %39 %80
%82 = OpAccessChain %79 %31 %81
%83 = OpLoad %10 %82
%84 = OpCompositeConstruct %13 %83 %83 %83 %83
%85 = OpFAdd %13 %78 %84
%88 = OpExtInst %45 %1 UMin %39 %87
%89 = OpAccessChain %86 %33 %88
%90 = OpLoad %10 %89
%91 = OpCompositeConstruct %13 %90 %90 %90 %90
%92 = OpFAdd %13 %85 %91
%95 = OpExtInst %45 %1 UMin %39 %94
%96 = OpAccessChain %93 %35 %95
%97 = OpLoad %13 %96
%98 = OpFAdd %13 %92 %97
OpReturnValue %98
OpFunctionEnd

276
tests/out/spv/shadow.spvasm
View File

@ -1,13 +1,13 @@
; SPIR-V
; Version: 1.2
; Generator: rspirv
; Bound: 122
; Bound: 125
OpCapability Shader
OpExtension "SPV_KHR_storage_buffer_storage_class"
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %77 "fs_main" %69 %72 %75
OpExecutionMode %77 OriginUpperLeft
OpEntryPoint Fragment %79 "fs_main" %71 %74 %77
OpExecutionMode %79 OriginUpperLeft
OpSource GLSL 450
OpName %9 "c_max_lights"
OpMemberName %14 0 "num_lights"
@ -20,18 +20,17 @@ OpMemberName %19 0 "data"
OpName %19 "Lights"
OpName %24 "c_ambient"
OpName %25 "u_globals"
OpName %27 "s_lights"
OpName %29 "t_shadow"
OpName %31 "sampler_shadow"
OpName %34 "light_id"
OpName %35 "homogeneous_coords"
OpName %36 "fetch_shadow"
OpName %64 "color"
OpName %66 "i"
OpName %69 "raw_normal"
OpName %72 "position"
OpName %77 "fs_main"
OpDecorate %14 Block
OpName %28 "s_lights"
OpName %30 "t_shadow"
OpName %32 "sampler_shadow"
OpName %35 "light_id"
OpName %36 "homogeneous_coords"
OpName %37 "fetch_shadow"
OpName %66 "color"
OpName %68 "i"
OpName %71 "raw_normal"
OpName %74 "position"
OpName %79 "fs_main"
OpMemberDecorate %14 0 Offset 0
OpMemberDecorate %17 0 Offset 0
OpMemberDecorate %17 0 ColMajor
@ -39,20 +38,22 @@ OpMemberDecorate %17 0 MatrixStride 16
OpMemberDecorate %17 1 Offset 64
OpMemberDecorate %17 2 Offset 80
OpDecorate %18 ArrayStride 96
OpDecorate %19 Block
OpMemberDecorate %19 0 Offset 0
OpDecorate %25 DescriptorSet 0
OpDecorate %25 Binding 0
OpDecorate %27 NonWritable
OpDecorate %27 DescriptorSet 0
OpDecorate %27 Binding 1
OpDecorate %29 DescriptorSet 0
OpDecorate %29 Binding 2
OpDecorate %31 DescriptorSet 0
OpDecorate %31 Binding 3
OpDecorate %69 Location 0
OpDecorate %72 Location 1
OpDecorate %75 Location 0
OpDecorate %26 Block
OpMemberDecorate %26 0 Offset 0
OpDecorate %28 NonWritable
OpDecorate %28 DescriptorSet 0
OpDecorate %28 Binding 1
OpDecorate %19 Block
OpDecorate %30 DescriptorSet 0
OpDecorate %30 Binding 2
OpDecorate %32 DescriptorSet 0
OpDecorate %32 Binding 3
OpDecorate %71 Location 0
OpDecorate %74 Location 1
OpDecorate %77 Location 0
%2 = OpTypeVoid
%4 = OpTypeFloat 32
%3 = OpConstant %4 0.0
@ -76,120 +77,123 @@ OpDecorate %75 Location 0
%22 = OpTypeVector %4 2
%23 = OpTypeVector %4 3
%24 = OpConstantComposite %23 %8 %8 %8
%26 = OpTypePointer Uniform %14
%25 = OpVariable %26 Uniform
%28 = OpTypePointer StorageBuffer %19
%27 = OpVariable %28 StorageBuffer
%30 = OpTypePointer UniformConstant %20
%29 = OpVariable %30 UniformConstant
%32 = OpTypePointer UniformConstant %21
%31 = OpVariable %32 UniformConstant
%37 = OpTypeFunction %4 %10 %16
%42 = OpTypeBool
%54 = OpTypeInt 32 1
%59 = OpTypeSampledImage %20
%65 = OpTypePointer Function %23
%67 = OpTypePointer Function %10
%70 = OpTypePointer Input %23
%69 = OpVariable %70 Input
%73 = OpTypePointer Input %16
%72 = OpVariable %73 Input
%76 = OpTypePointer Output %16
%75 = OpVariable %76 Output
%78 = OpTypeFunction %2
%88 = OpTypePointer Uniform %13
%89 = OpTypePointer Uniform %10
%96 = OpTypePointer StorageBuffer %18
%98 = OpTypePointer StorageBuffer %17
%36 = OpFunction %4 None %37
%34 = OpFunctionParameter %10
%35 = OpFunctionParameter %16
%33 = OpLabel
%38 = OpLoad %20 %29
%39 = OpLoad %21 %31
OpBranch %40
%40 = OpLabel
%41 = OpCompositeExtract %4 %35 3
%43 = OpFOrdLessThanEqual %42 %41 %3
OpSelectionMerge %44 None
OpBranchConditional %43 %45 %44
%45 = OpLabel
%26 = OpTypeStruct %14
%27 = OpTypePointer Uniform %26
%25 = OpVariable %27 Uniform
%29 = OpTypePointer StorageBuffer %19
%28 = OpVariable %29 StorageBuffer
%31 = OpTypePointer UniformConstant %20
%30 = OpVariable %31 UniformConstant
%33 = OpTypePointer UniformConstant %21
%32 = OpVariable %33 UniformConstant
%38 = OpTypeFunction %4 %10 %16
%41 = OpTypePointer Uniform %14
%44 = OpTypeBool
%56 = OpTypeInt 32 1
%61 = OpTypeSampledImage %20
%67 = OpTypePointer Function %23
%69 = OpTypePointer Function %10
%72 = OpTypePointer Input %23
%71 = OpVariable %72 Input
%75 = OpTypePointer Input %16
%74 = OpVariable %75 Input
%78 = OpTypePointer Output %16
%77 = OpVariable %78 Output
%80 = OpTypeFunction %2
%91 = OpTypePointer Uniform %13
%92 = OpTypePointer Uniform %10
%99 = OpTypePointer StorageBuffer %18
%101 = OpTypePointer StorageBuffer %17
%37 = OpFunction %4 None %38
%35 = OpFunctionParameter %10
%36 = OpFunctionParameter %16
%34 = OpLabel
%39 = OpLoad %20 %30
%40 = OpLoad %21 %32
OpBranch %42
%42 = OpLabel
%43 = OpCompositeExtract %4 %36 3
%45 = OpFOrdLessThanEqual %44 %43 %3
OpSelectionMerge %46 None
OpBranchConditional %45 %47 %46
%47 = OpLabel
OpReturnValue %5
%44 = OpLabel
%46 = OpCompositeConstruct %22 %6 %7
%47 = OpVectorShuffle %22 %35 %35 0 1
%48 = OpFMul %22 %47 %46
%49 = OpCompositeExtract %4 %35 3
%50 = OpCompositeConstruct %22 %49 %49
%51 = OpFDiv %22 %48 %50
%52 = OpCompositeConstruct %22 %6 %6
%53 = OpFAdd %22 %51 %52
%55 = OpBitcast %54 %34
%56 = OpCompositeExtract %4 %35 2
%57 = OpCompositeExtract %4 %35 3
%58 = OpFDiv %4 %56 %57
%60 = OpConvertUToF %4 %55
%61 = OpCompositeConstruct %23 %53 %60
%62 = OpSampledImage %59 %38 %39
%63 = OpImageSampleDrefExplicitLod %4 %62 %61 %58 Lod %3
OpReturnValue %63
%46 = OpLabel
%48 = OpCompositeConstruct %22 %6 %7
%49 = OpVectorShuffle %22 %36 %36 0 1
%50 = OpFMul %22 %49 %48
%51 = OpCompositeExtract %4 %36 3
%52 = OpCompositeConstruct %22 %51 %51
%53 = OpFDiv %22 %50 %52
%54 = OpCompositeConstruct %22 %6 %6
%55 = OpFAdd %22 %53 %54
%57 = OpBitcast %56 %35
%58 = OpCompositeExtract %4 %36 2
%59 = OpCompositeExtract %4 %36 3
%60 = OpFDiv %4 %58 %59
%62 = OpConvertUToF %4 %57
%63 = OpCompositeConstruct %23 %55 %62
%64 = OpSampledImage %61 %39 %40
%65 = OpImageSampleDrefExplicitLod %4 %64 %63 %60 Lod %3
OpReturnValue %65
OpFunctionEnd
%77 = OpFunction %2 None %78
%68 = OpLabel
%64 = OpVariable %65 Function %24
%66 = OpVariable %67 Function %11
%71 = OpLoad %23 %69
%74 = OpLoad %16 %72
%79 = OpLoad %20 %29
%80 = OpLoad %21 %31
OpBranch %81
%81 = OpLabel
%82 = OpExtInst %23 %1 Normalize %71
OpBranch %83
%83 = OpLabel
OpLoopMerge %84 %86 None
OpBranch %85
%85 = OpLabel
%87 = OpLoad %10 %66
%90 = OpAccessChain %89 %25 %11 %11
%91 = OpLoad %10 %90
%92 = OpExtInst %10 %1 UMin %91 %9
%93 = OpUGreaterThanEqual %42 %87 %92
OpSelectionMerge %94 None
OpBranchConditional %93 %95 %94
%95 = OpLabel
%79 = OpFunction %2 None %80
%70 = OpLabel
%66 = OpVariable %67 Function %24
%68 = OpVariable %69 Function %11
%73 = OpLoad %23 %71
%76 = OpLoad %16 %74
%81 = OpAccessChain %41 %25 %11
%82 = OpLoad %20 %30
%83 = OpLoad %21 %32
OpBranch %84
%94 = OpLabel
%97 = OpLoad %10 %66
%99 = OpAccessChain %98 %27 %11 %97
%100 = OpLoad %17 %99
%101 = OpLoad %10 %66
%102 = OpCompositeExtract %15 %100 0
%103 = OpMatrixTimesVector %16 %102 %74
%104 = OpFunctionCall %4 %36 %101 %103
%105 = OpCompositeExtract %16 %100 1
%106 = OpVectorShuffle %23 %105 %105 0 1 2
%107 = OpVectorShuffle %23 %74 %74 0 1 2
%108 = OpFSub %23 %106 %107
%109 = OpExtInst %23 %1 Normalize %108
%110 = OpDot %4 %82 %109
%111 = OpExtInst %4 %1 FMax %3 %110
%112 = OpLoad %23 %64
%113 = OpFMul %4 %104 %111
%114 = OpCompositeExtract %16 %100 2
%115 = OpVectorShuffle %23 %114 %114 0 1 2
%116 = OpVectorTimesScalar %23 %115 %113
%117 = OpFAdd %23 %112 %116
OpStore %64 %117
%84 = OpLabel
%85 = OpExtInst %23 %1 Normalize %73
OpBranch %86
%86 = OpLabel
%118 = OpLoad %10 %66
%119 = OpIAdd %10 %118 %12
OpStore %66 %119
OpBranch %83
%84 = OpLabel
%120 = OpLoad %23 %64
%121 = OpCompositeConstruct %16 %120 %5
OpStore %75 %121
OpLoopMerge %87 %89 None
OpBranch %88
%88 = OpLabel
%90 = OpLoad %10 %68
%93 = OpAccessChain %92 %81 %11 %11
%94 = OpLoad %10 %93
%95 = OpExtInst %10 %1 UMin %94 %9
%96 = OpUGreaterThanEqual %44 %90 %95
OpSelectionMerge %97 None
OpBranchConditional %96 %98 %97
%98 = OpLabel
OpBranch %87
%97 = OpLabel
%100 = OpLoad %10 %68
%102 = OpAccessChain %101 %28 %11 %100
%103 = OpLoad %17 %102
%104 = OpLoad %10 %68
%105 = OpCompositeExtract %15 %103 0
%106 = OpMatrixTimesVector %16 %105 %76
%107 = OpFunctionCall %4 %37 %104 %106
%108 = OpCompositeExtract %16 %103 1
%109 = OpVectorShuffle %23 %108 %108 0 1 2
%110 = OpVectorShuffle %23 %76 %76 0 1 2
%111 = OpFSub %23 %109 %110
%112 = OpExtInst %23 %1 Normalize %111
%113 = OpDot %4 %85 %112
%114 = OpExtInst %4 %1 FMax %3 %113
%115 = OpLoad %23 %66
%116 = OpFMul %4 %107 %114
%117 = OpCompositeExtract %16 %103 2
%118 = OpVectorShuffle %23 %117 %117 0 1 2
%119 = OpVectorTimesScalar %23 %118 %116
%120 = OpFAdd %23 %115 %119
OpStore %66 %120
OpBranch %89
%89 = OpLabel
%121 = OpLoad %10 %68
%122 = OpIAdd %10 %121 %12
OpStore %68 %122
OpBranch %86
%87 = OpLabel
%123 = OpLoad %23 %66
%124 = OpCompositeConstruct %16 %123 %5
OpStore %77 %124
OpReturn
OpFunctionEnd

182
tests/out/spv/skybox.spvasm
View File

@ -1,16 +1,15 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 93
; Bound: 96
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %36 "vs_main" %29 %32 %34
OpEntryPoint Fragment %85 "fs_main" %78 %81 %84
OpExecutionMode %85 OriginUpperLeft
OpEntryPoint Vertex %37 "vs_main" %30 %33 %35
OpEntryPoint Fragment %88 "fs_main" %81 %84 %87
OpExecutionMode %88 OriginUpperLeft
OpMemberDecorate %12 0 Offset 0
OpMemberDecorate %12 1 Offset 16
OpDecorate %14 Block
OpMemberDecorate %14 0 Offset 0
OpMemberDecorate %14 0 ColMajor
OpMemberDecorate %14 0 MatrixStride 16
@ -19,16 +18,18 @@ OpMemberDecorate %14 1 ColMajor
OpMemberDecorate %14 1 MatrixStride 16
OpDecorate %19 DescriptorSet 0
OpDecorate %19 Binding 0
OpDecorate %21 DescriptorSet 0
OpDecorate %21 Binding 1
OpDecorate %23 DescriptorSet 0
OpDecorate %23 Binding 2
OpDecorate %29 BuiltIn VertexIndex
OpDecorate %32 BuiltIn Position
OpDecorate %34 Location 0
OpDecorate %78 BuiltIn FragCoord
OpDecorate %81 Location 0
OpDecorate %20 Block
OpMemberDecorate %20 0 Offset 0
OpDecorate %22 DescriptorSet 0
OpDecorate %22 Binding 1
OpDecorate %24 DescriptorSet 0
OpDecorate %24 Binding 2
OpDecorate %30 BuiltIn VertexIndex
OpDecorate %33 BuiltIn Position
OpDecorate %35 Location 0
OpDecorate %81 BuiltIn FragCoord
OpDecorate %84 Location 0
OpDecorate %87 Location 0
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpConstant %4 2
@ -46,88 +47,91 @@ OpDecorate %84 Location 0
%16 = OpTypeMatrix %11 3
%17 = OpTypeImage %7 Cube 0 0 0 1 Unknown
%18 = OpTypeSampler
%20 = OpTypePointer Uniform %14
%19 = OpVariable %20 Uniform
%22 = OpTypePointer UniformConstant %17
%21 = OpVariable %22 UniformConstant
%24 = OpTypePointer UniformConstant %18
%23 = OpVariable %24 UniformConstant
%26 = OpTypePointer Function %4
%30 = OpTypePointer Input %15
%29 = OpVariable %30 Input
%33 = OpTypePointer Output %10
%32 = OpVariable %33 Output
%35 = OpTypePointer Output %11
%34 = OpVariable %35 Output
%37 = OpTypeFunction %2
%52 = OpTypePointer Uniform %13
%53 = OpTypePointer Uniform %10
%54 = OpConstant %15 0
%55 = OpConstant %15 1
%62 = OpConstant %15 2
%79 = OpTypePointer Input %10
%78 = OpVariable %79 Input
%82 = OpTypePointer Input %11
%20 = OpTypeStruct %14
%21 = OpTypePointer Uniform %20
%19 = OpVariable %21 Uniform
%23 = OpTypePointer UniformConstant %17
%22 = OpVariable %23 UniformConstant
%25 = OpTypePointer UniformConstant %18
%24 = OpVariable %25 UniformConstant
%27 = OpTypePointer Function %4
%31 = OpTypePointer Input %15
%30 = OpVariable %31 Input
%34 = OpTypePointer Output %10
%33 = OpVariable %34 Output
%36 = OpTypePointer Output %11
%35 = OpVariable %36 Output
%38 = OpTypeFunction %2
%39 = OpTypePointer Uniform %14
%40 = OpConstant %15 0
%56 = OpTypePointer Uniform %13
%57 = OpTypePointer Uniform %10
%58 = OpConstant %15 1
%65 = OpConstant %15 2
%82 = OpTypePointer Input %10
%81 = OpVariable %82 Input
%84 = OpVariable %33 Output
%90 = OpTypeSampledImage %17
%36 = OpFunction %2 None %37
%28 = OpLabel
%25 = OpVariable %26 Function
%27 = OpVariable %26 Function
%31 = OpLoad %15 %29
OpBranch %38
%38 = OpLabel
%39 = OpBitcast %4 %31
%40 = OpSDiv %4 %39 %3
OpStore %25 %40
%41 = OpBitcast %4 %31
%42 = OpBitwiseAnd %4 %41 %5
OpStore %27 %42
%43 = OpLoad %4 %25
%44 = OpConvertSToF %7 %43
%45 = OpFMul %7 %44 %6
%46 = OpFSub %7 %45 %8
%47 = OpLoad %4 %27
%85 = OpTypePointer Input %11
%84 = OpVariable %85 Input
%87 = OpVariable %34 Output
%93 = OpTypeSampledImage %17
%37 = OpFunction %2 None %38
%29 = OpLabel
%26 = OpVariable %27 Function
%28 = OpVariable %27 Function
%32 = OpLoad %15 %30
%41 = OpAccessChain %39 %19 %40
OpBranch %42
%42 = OpLabel
%43 = OpBitcast %4 %32
%44 = OpSDiv %4 %43 %3
OpStore %26 %44
%45 = OpBitcast %4 %32
%46 = OpBitwiseAnd %4 %45 %5
OpStore %28 %46
%47 = OpLoad %4 %26
%48 = OpConvertSToF %7 %47
%49 = OpFMul %7 %48 %6
%50 = OpFSub %7 %49 %8
%51 = OpCompositeConstruct %10 %46 %50 %9 %8
%56 = OpAccessChain %53 %19 %55 %54
%57 = OpLoad %10 %56
%58 = OpVectorShuffle %11 %57 %57 0 1 2
%59 = OpAccessChain %53 %19 %55 %55
%51 = OpLoad %4 %28
%52 = OpConvertSToF %7 %51
%53 = OpFMul %7 %52 %6
%54 = OpFSub %7 %53 %8
%55 = OpCompositeConstruct %10 %50 %54 %9 %8
%59 = OpAccessChain %57 %41 %58 %40
%60 = OpLoad %10 %59
%61 = OpVectorShuffle %11 %60 %60 0 1 2
%63 = OpAccessChain %53 %19 %55 %62
%64 = OpLoad %10 %63
%65 = OpVectorShuffle %11 %64 %64 0 1 2
%66 = OpCompositeConstruct %16 %58 %61 %65
%67 = OpTranspose %16 %66
%68 = OpAccessChain %52 %19 %54
%69 = OpLoad %13 %68
%70 = OpMatrixTimesVector %10 %69 %51
%71 = OpVectorShuffle %11 %70 %70 0 1 2
%72 = OpMatrixTimesVector %11 %67 %71
%73 = OpCompositeConstruct %12 %51 %72
%74 = OpCompositeExtract %10 %73 0
OpStore %32 %74
%75 = OpCompositeExtract %11 %73 1
OpStore %34 %75
%62 = OpAccessChain %57 %41 %58 %58
%63 = OpLoad %10 %62
%64 = OpVectorShuffle %11 %63 %63 0 1 2
%66 = OpAccessChain %57 %41 %58 %65
%67 = OpLoad %10 %66
%68 = OpVectorShuffle %11 %67 %67 0 1 2
%69 = OpCompositeConstruct %16 %61 %64 %68
%70 = OpTranspose %16 %69
%71 = OpAccessChain %56 %41 %40
%72 = OpLoad %13 %71
%73 = OpMatrixTimesVector %10 %72 %55
%74 = OpVectorShuffle %11 %73 %73 0 1 2
%75 = OpMatrixTimesVector %11 %70 %74
%76 = OpCompositeConstruct %12 %55 %75
%77 = OpCompositeExtract %10 %76 0
OpStore %33 %77
%78 = OpCompositeExtract %11 %76 1
OpStore %35 %78
OpReturn
OpFunctionEnd
%85 = OpFunction %2 None %37
%76 = OpLabel
%80 = OpLoad %10 %78
%83 = OpLoad %11 %81
%77 = OpCompositeConstruct %12 %80 %83
%86 = OpLoad %17 %21
%87 = OpLoad %18 %23
OpBranch %88
%88 = OpLabel
%89 = OpCompositeExtract %11 %77 1
%91 = OpSampledImage %90 %86 %87
%92 = OpImageSampleImplicitLod %10 %91 %89
OpStore %84 %92
%88 = OpFunction %2 None %38
%79 = OpLabel
%83 = OpLoad %10 %81
%86 = OpLoad %11 %84
%80 = OpCompositeConstruct %12 %83 %86
%89 = OpLoad %17 %22
%90 = OpLoad %18 %24
OpBranch %91
%91 = OpLabel
%92 = OpCompositeExtract %11 %80 1
%94 = OpSampledImage %93 %89 %90
%95 = OpImageSampleImplicitLod %10 %94 %92
OpStore %87 %95
OpReturn
OpFunctionEnd

1
tests/out/wgsl/210-bevy-2d-shader-frag.wgsl
View File

@ -1,4 +1,3 @@
[[block]]
struct ColorMaterial_color {
Color: vec4<f32>;
};

3
tests/out/wgsl/210-bevy-2d-shader-vert.wgsl
View File

@ -1,14 +1,11 @@
[[block]]
struct Camera {
ViewProj: mat4x4<f32>;
};
[[block]]
struct Transform {
Model: mat4x4<f32>;
};
[[block]]
struct Sprite_size {
size: vec2<f32>;
};

2
tests/out/wgsl/210-bevy-shader-vert.wgsl
View File

@ -1,9 +1,7 @@
[[block]]
struct Camera {
ViewProj: mat4x4<f32>;
};
[[block]]
struct Transform {
Model: mat4x4<f32>;
};

1
tests/out/wgsl/246-collatz-comp.wgsl
View File

@ -1,4 +1,3 @@
[[block]]
struct PrimeIndices {
indices: [[stride(4)]] array<u32>;
};

2
tests/out/wgsl/800-out-of-bounds-panic-vert.wgsl
View File

@ -1,9 +1,7 @@
[[block]]
struct Globals {
view_matrix: mat4x4<f32>;
};
[[block]]
struct VertexPushConstants {
world_matrix: mat4x4<f32>;
};

1
tests/out/wgsl/896-push-constant-vert.wgsl
View File

@ -1,4 +1,3 @@
[[block]]
struct PushConstants {
example: f32;
};

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

@ -1,4 +1,3 @@
[[block]]
struct Bar {
matrix: mat4x4<f32>;
atom: atomic<i32>;

8
tests/out/wgsl/bevy-pbr-frag.wgsl
View File

@ -9,17 +9,14 @@ struct DirectionalLight {
color: vec4<f32>;
};
[[block]]
struct CameraViewProj {
ViewProj: mat4x4<f32>;
};
[[block]]
struct CameraPosition {
CameraPos: vec4<f32>;
};
[[block]]
struct Lights {
AmbientColor: vec4<f32>;
NumLights: vec4<u32>;
@ -27,27 +24,22 @@ struct Lights {
DirectionalLights: [[stride(32)]] array<DirectionalLight,1u>;
};
[[block]]
struct StandardMaterial_base_color {
base_color: vec4<f32>;
};
[[block]]
struct StandardMaterial_roughness {
perceptual_roughness: f32;
};
[[block]]
struct StandardMaterial_metallic {
metallic: f32;
};
[[block]]
struct StandardMaterial_reflectance {
reflectance: f32;
};
[[block]]
struct StandardMaterial_emissive {
emissive: vec4<f32>;
};

2
tests/out/wgsl/bevy-pbr-vert.wgsl
View File

@ -1,9 +1,7 @@
[[block]]
struct CameraViewProj {
ViewProj: mat4x4<f32>;
};
[[block]]
struct Transform {
Model: mat4x4<f32>;
};

2
tests/out/wgsl/boids.wgsl
View File

@ -3,7 +3,6 @@ struct Particle {
vel: vec2<f32>;
};
[[block]]
struct SimParams {
deltaT: f32;
rule1Distance: f32;
@ -14,7 +13,6 @@ struct SimParams {
rule3Scale: f32;
};
[[block]]
struct Particles {
particles: [[stride(16)]] array<Particle>;
};

1
tests/out/wgsl/collatz.wgsl
View File

@ -1,4 +1,3 @@
[[block]]
struct PrimeIndices {
data: [[stride(4)]] array<u32>;
};

1
tests/out/wgsl/constant-array-size-vert.wgsl
View File

@ -1,4 +1,3 @@
[[block]]
struct Data {
vecs: [[stride(16)]] array<vec4<f32>,42u>;
};

1
tests/out/wgsl/empty-global-name.wgsl
View File

@ -1,4 +1,3 @@
[[block]]
struct type_1 {
member: i32;
};

1
tests/out/wgsl/extra.wgsl
View File

@ -1,4 +1,3 @@
[[block]]
struct PushConstants {
index: u32;
double: vec2<f32>;

1
tests/out/wgsl/pointers.wgsl
View File

@ -1,4 +1,3 @@
[[block]]
struct DynamicArray {
arr: [[stride(4)]] array<u32>;
};

1
tests/out/wgsl/quad-vert.wgsl
View File

@ -1,4 +1,3 @@
[[block]]
struct gl_PerVertex {
[[builtin(position)]] gl_Position: vec4<f32>;
};

2
tests/out/wgsl/shadow.wgsl
View File

@ -1,4 +1,3 @@
[[block]]
struct Globals {
num_lights: vec4<u32>;
};
@ -9,7 +8,6 @@ struct Light {
color: vec4<f32>;
};
[[block]]
struct Lights {
data: [[stride(96)]] array<Light>;
};

1
tests/out/wgsl/skybox.wgsl
View File

@ -3,7 +3,6 @@ struct VertexOutput {
[[location(0)]] uv: vec3<f32>;
};
[[block]]
struct Data {
proj_inv: mat4x4<f32>;
view: mat4x4<f32>;

14
tests/wgsl-errors.rs
View File

@ -804,17 +804,6 @@ fn invalid_arrays() {
})
}
check_validation_error! {
r#"
[[block]] struct Block { value: f32; };
type Bad = array<Block, 4>;
"#:
Err(naga::valid::ValidationError::Type {
error: naga::valid::TypeError::NestedTopLevel,
..
})
}
check_validation_error! {
r#"
type Bad = [[stride(2)]] array<f32, 4>;
@ -1154,7 +1143,6 @@ fn invalid_runtime_sized_arrays() {
arr: array<f32>;
};
[[block]]
struct Outer {
legit: i32;
unsized: Unsized;
@ -1264,7 +1252,6 @@ fn wrong_access_mode() {
// variables whose access mode is `read`, not `read_write`.
check_validation_error! {
"
[[block]]
struct Globals {
i: i32;
};
@ -1277,7 +1264,6 @@ fn wrong_access_mode() {
}
",
"
[[block]]
struct Globals {
i: i32;
};