From 8ffd6ba929b4b93c9564f08fe8bb34b23fa72a6f Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Wed, 8 Dec 2021 00:27:20 -0500 Subject: [PATCH] Remove top_level property of structs --- CHANGELOG.md | 6 +- src/back/glsl/mod.rs | 146 +++-- src/back/hlsl/writer.rs | 15 +- src/back/msl/writer.rs | 4 +- src/back/spv/block.rs | 4 +- src/back/spv/helpers.rs | 26 + src/back/spv/index.rs | 7 +- src/back/spv/mod.rs | 27 +- src/back/spv/writer.rs | 107 +++- src/back/wgsl/writer.rs | 13 +- src/front/glsl/functions.rs | 6 +- src/front/glsl/offset.rs | 12 +- src/front/glsl/parser/declarations.rs | 1 - src/front/glsl/parser/types.rs | 6 +- src/front/spv/function.rs | 1 - src/front/spv/mod.rs | 24 +- src/front/wgsl/mod.rs | 12 +- src/front/wgsl/tests.rs | 3 +- src/lib.rs | 14 +- src/proc/layouter.rs | 6 +- src/valid/interface.rs | 17 +- src/valid/type.rs | 20 +- tests/in/access.wgsl | 1 - tests/in/boids.wgsl | 2 - tests/in/bounds-check-restrict.wgsl | 1 - tests/in/bounds-check-zero.wgsl | 1 - tests/in/collatz.wgsl | 1 - tests/in/extra.wgsl | 1 - tests/in/pointers.wgsl | 1 - tests/in/policy-mix.wgsl | 2 - tests/in/shadow.wgsl | 2 - tests/in/skybox.wgsl | 1 - tests/out/glsl/access.atomics.Compute.glsl | 2 +- tests/out/glsl/access.foo.Vertex.glsl | 2 +- tests/out/glsl/boids.main.Compute.glsl | 10 +- .../glsl/interpolate.frag_main.Fragment.glsl | 1 - tests/out/glsl/interpolate.main.Fragment.glsl | 25 - tests/out/glsl/interpolate.main.Vertex.glsl | 42 -- .../glsl/interpolate.vert_main.Vertex.glsl | 1 - tests/out/glsl/operators.main.Compute.glsl | 1 - tests/out/glsl/quad-vert.main.Vertex.glsl | 14 +- tests/out/glsl/quad.frag_main.Fragment.glsl | 1 - tests/out/glsl/quad.fs_extra.Fragment.glsl | 1 - tests/out/glsl/quad.vert_main.Vertex.glsl | 1 - tests/out/glsl/shadow.fs_main.Fragment.glsl | 10 +- tests/out/glsl/skybox.fs_main.Fragment.glsl | 5 +- tests/out/glsl/skybox.vs_main.Vertex.glsl | 6 +- tests/out/ir/collatz.ron | 1 - tests/out/ir/shadow.ron | 3 - tests/out/spv/access.spvasm | 2 +- tests/out/spv/boids.spvasm | 539 +++++++++--------- tests/out/spv/bounds-check-restrict.spvasm | 2 +- tests/out/spv/bounds-check-zero.spvasm | 2 +- tests/out/spv/collatz.spvasm | 2 +- tests/out/spv/extra.spvasm | 1 - tests/out/spv/pointers.spvasm | 2 +- tests/out/spv/policy-mix.spvasm | 192 ++++--- tests/out/spv/shadow.spvasm | 276 ++++----- tests/out/spv/skybox.spvasm | 182 +++--- tests/out/wgsl/210-bevy-2d-shader-frag.wgsl | 1 - tests/out/wgsl/210-bevy-2d-shader-vert.wgsl | 3 - tests/out/wgsl/210-bevy-shader-vert.wgsl | 2 - tests/out/wgsl/246-collatz-comp.wgsl | 1 - .../wgsl/800-out-of-bounds-panic-vert.wgsl | 2 - tests/out/wgsl/896-push-constant-vert.wgsl | 1 - tests/out/wgsl/access.wgsl | 1 - tests/out/wgsl/bevy-pbr-frag.wgsl | 8 - tests/out/wgsl/bevy-pbr-vert.wgsl | 2 - tests/out/wgsl/boids.wgsl | 2 - tests/out/wgsl/collatz.wgsl | 1 - tests/out/wgsl/constant-array-size-vert.wgsl | 1 - tests/out/wgsl/empty-global-name.wgsl | 1 - tests/out/wgsl/extra.wgsl | 1 - tests/out/wgsl/pointers.wgsl | 1 - tests/out/wgsl/quad-vert.wgsl | 1 - tests/out/wgsl/shadow.wgsl | 2 - tests/out/wgsl/skybox.wgsl | 1 - tests/wgsl-errors.rs | 14 - 78 files changed, 882 insertions(+), 980 deletions(-) delete mode 100644 tests/out/glsl/interpolate.main.Fragment.glsl delete mode 100644 tests/out/glsl/interpolate.main.Vertex.glsl diff --git a/CHANGELOG.md b/CHANGELOG.md index 8439ecbc74..da938ff2e1 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -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) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 53f376a660..082fa685c5 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -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) -> 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, - /// A map with all the names needed for reflections - reflection_names_uniforms: crate::FastHashMap, String>, /// A map with the names of global variables needed for reflections reflection_names_globals: crate::FastHashMap, 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); - - if !is_global_struct { - self.write_struct(false, handle, members)? + let used_by_global = self.module.global_variables.iter().any(|(vh, var)| { + !ep_info[vh].is_empty() && var.class.is_buffer() && var.ty == handle + }); + + 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, 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); } _ => (), diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 6f05f1bd10..d8cfde3cd0 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -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, - _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, diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index 3a917a3562..51f892d783 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -2652,9 +2652,7 @@ impl Writer { 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, diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index f751cba508..f0676f0d40 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -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]; diff --git a/src/back/spv/helpers.rs b/src/back/spv/helpers.rs index aa7952ad29..4640ebf524 100644 --- a/src/back/spv/helpers.rs +++ b/src/back/spv/helpers.rs @@ -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, + } +} diff --git a/src/back/spv/index.rs b/src/back/spv/index.rs index 4192d3848e..c2b198048a 100644 --- a/src/back/spv/index.rs +++ b/src/back/spv/index.rs @@ -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")), } } diff --git a/src/back/spv/mod.rs b/src/back/spv/mod.rs index d341e8df2d..8dc574dfe8 100644 --- a/src/back/spv/mod.rs +++ b/src/back/spv/mod.rs @@ -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; } } diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs index 73066c9055..ea83a42031 100644 --- a/src/back/spv/writer.rs +++ b/src/back/spv/writer.rs @@ -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> { - let id = self.id_gen.next(); + ) -> Result { + 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) } }; diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index 0f19f765b2..4d06b7a941 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -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, Option), @@ -107,12 +106,11 @@ impl Writer { // 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 Writer { 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 Writer { &mut self, module: &Module, handle: Handle, - 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, " {{")?; diff --git a/src/front/glsl/functions.rs b/src/front/glsl/functions.rs index acaa5c9759..0bc44b1015 100644 --- a/src/front/glsl/functions.rs +++ b/src/front/glsl/functions.rs @@ -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(), ); diff --git a/src/front/glsl/offset.rs b/src/front/glsl/offset.rs index 7a7b5e2f6e..dd8ecbfea9 100644 --- a/src/front/glsl/offset.rs +++ b/src/front/glsl/offset.rs @@ -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, ); diff --git a/src/front/glsl/parser/declarations.rs b/src/front/glsl/parser/declarations.rs index 2cb4f43e0f..4911d0734d 100644 --- a/src/front/glsl/parser/declarations.rs +++ b/src/front/glsl/parser/declarations.rs @@ -501,7 +501,6 @@ impl<'source> ParsingContext<'source> { Type { name: Some(ty_name), inner: TypeInner::Struct { - top_level: true, members: members.clone(), span, }, diff --git a/src/front/glsl/parser/types.rs b/src/front/glsl/parser/types.rs index 7fc74e2ddf..ad537e9594 100644 --- a/src/front/glsl/parser/types.rs +++ b/src/front/glsl/parser/types.rs @@ -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, ); diff --git a/src/front/spv/function.rs b/src/front/spv/function.rs index d85c61438c..23c47688e7 100644 --- a/src/front/spv/function.rs +++ b/src/front/spv/function.rs @@ -488,7 +488,6 @@ impl> super::Parser { crate::Type { name: None, inner: crate::TypeInner::Struct { - top_level: false, members, span: 0xFFFF, // shouldn't matter }, diff --git a/src/front/spv/mod.rs b/src/front/spv/mod.rs index 55a5691dee..cd1721bbd5 100644 --- a/src/front/spv/mod.rs +++ b/src/front/spv/mod.rs @@ -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, desc_index: Option, specialization: Option, - block: Option, + storage_buffer: bool, offset: Option, array_stride: Option, matrix_stride: Option, @@ -663,11 +658,8 @@ impl> Parser { 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> Parser { 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> Parser { 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> Parser { 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); } diff --git a/src/front/wgsl/mod.rs b/src/front/wgsl/mod.rs index 9dfc194f2f..0e39ee176b 100644 --- a/src/front/wgsl/mod.rs +++ b/src/front/wgsl/mod.rs @@ -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, ); diff --git a/src/front/wgsl/tests.rs b/src/front/wgsl/tests.rs index 1434320633..f64ac30369 100644 --- a/src/front/wgsl/tests.rs +++ b/src/front/wgsl/tests.rs @@ -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; [[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; }; // this is used as both input and output for convenience diff --git a/src/lib.rs b/src/lib.rs index eea9d2e932..a2d2ec15a7 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -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, //TODO: should this be unaligned? span: u32, diff --git a/src/proc/layouter.rs b/src/proc/layouter.rs index 5facecdac2..5f6559f08e 100644 --- a/src/proc/layouter.rs +++ b/src/proc/layouter.rs @@ -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 { diff --git a/src/valid/interface.rs b/src/valid/interface.rs index 68326b8a8d..1fabe5f3b0 100644 --- a/src/valid/interface.rs +++ b/src/valid/interface.rs @@ -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, ) } diff --git a/src/valid/type.rs b/src/valid/type.rs index 923e1b24ff..cda1fa64a0 100644 --- a/src/valid/type.rs +++ b/src/valid/type.rs @@ -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 { diff --git a/tests/in/access.wgsl b/tests/in/access.wgsl index 034d0d5ab2..e8ef082fbc 100644 --- a/tests/in/access.wgsl +++ b/tests/in/access.wgsl @@ -1,6 +1,5 @@ // This snapshot tests accessing various containers, dereferencing pointers. -[[block]] struct Bar { matrix: mat4x4; atom: atomic; diff --git a/tests/in/boids.wgsl b/tests/in/boids.wgsl index 0f89f9a203..5e712ac56c 100644 --- a/tests/in/boids.wgsl +++ b/tests/in/boids.wgsl @@ -5,7 +5,6 @@ struct Particle { vel : vec2; }; -[[block]] struct SimParams { deltaT : f32; rule1Distance : f32; @@ -16,7 +15,6 @@ struct SimParams { rule3Scale : f32; }; -[[block]] struct Particles { particles : [[stride(16)]] array; }; diff --git a/tests/in/bounds-check-restrict.wgsl b/tests/in/bounds-check-restrict.wgsl index 6f4ae998d8..8550fd7033 100644 --- a/tests/in/bounds-check-restrict.wgsl +++ b/tests/in/bounds-check-restrict.wgsl @@ -1,6 +1,5 @@ // Tests for `naga::back::BoundsCheckPolicy::Restrict`. -[[block]] struct Globals { a: array; v: vec4; diff --git a/tests/in/bounds-check-zero.wgsl b/tests/in/bounds-check-zero.wgsl index b4fbebcd78..4838701c23 100644 --- a/tests/in/bounds-check-zero.wgsl +++ b/tests/in/bounds-check-zero.wgsl @@ -1,6 +1,5 @@ // Tests for `naga::back::BoundsCheckPolicy::ReadZeroSkipWrite`. -[[block]] struct Globals { a: array; v: vec4; diff --git a/tests/in/collatz.wgsl b/tests/in/collatz.wgsl index 1db3b9115e..878b18d09b 100644 --- a/tests/in/collatz.wgsl +++ b/tests/in/collatz.wgsl @@ -1,4 +1,3 @@ -[[block]] struct PrimeIndices { data: [[stride(4)]] array; }; // this is used as both input and output for convenience diff --git a/tests/in/extra.wgsl b/tests/in/extra.wgsl index 9f9acb033c..fa8c466c4d 100644 --- a/tests/in/extra.wgsl +++ b/tests/in/extra.wgsl @@ -1,4 +1,3 @@ -[[block]] struct PushConstants { index: u32; double: vec2; diff --git a/tests/in/pointers.wgsl b/tests/in/pointers.wgsl index e3160774c3..cf365f3904 100644 --- a/tests/in/pointers.wgsl +++ b/tests/in/pointers.wgsl @@ -4,7 +4,6 @@ fn f() { *px = 10; } -[[block]] struct DynamicArray { arr: array; }; diff --git a/tests/in/policy-mix.wgsl b/tests/in/policy-mix.wgsl index e9b55709c9..4dc1e1bce8 100644 --- a/tests/in/policy-mix.wgsl +++ b/tests/in/policy-mix.wgsl @@ -2,13 +2,11 @@ // implemented separately. // Storage and Uniform storage classes -[[block]] struct InStorage { a: array, 10>; }; [[group(0), binding(0)]] var in_storage: InStorage; -[[block]] struct InUniform { a: array, 20>; }; diff --git a/tests/in/shadow.wgsl b/tests/in/shadow.wgsl index d1286273cc..f17f7af395 100644 --- a/tests/in/shadow.wgsl +++ b/tests/in/shadow.wgsl @@ -1,4 +1,3 @@ -[[block]] struct Globals { num_lights: vec4; }; @@ -12,7 +11,6 @@ struct Light { color: vec4; }; -[[block]] struct Lights { data: [[stride(96)]] array; }; diff --git a/tests/in/skybox.wgsl b/tests/in/skybox.wgsl index 33a1108406..b21624853f 100644 --- a/tests/in/skybox.wgsl +++ b/tests/in/skybox.wgsl @@ -3,7 +3,6 @@ struct VertexOutput { [[location(0)]] uv: vec3; }; -[[block]] struct Data { proj_inv: mat4x4; view: mat4x4; diff --git a/tests/out/glsl/access.atomics.Compute.glsl b/tests/out/glsl/access.atomics.Compute.glsl index 0145ad4981..19414a4496 100644 --- a/tests/out/glsl/access.atomics.Compute.glsl +++ b/tests/out/glsl/access.atomics.Compute.glsl @@ -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]; diff --git a/tests/out/glsl/access.foo.Vertex.glsl b/tests/out/glsl/access.foo.Vertex.glsl index d4502b0313..dfdf2c0130 100644 --- a/tests/out/glsl/access.foo.Vertex.glsl +++ b/tests/out/glsl/access.foo.Vertex.glsl @@ -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]; diff --git a/tests/out/glsl/boids.main.Compute.glsl b/tests/out/glsl/boids.main.Compute.glsl index 51bf0101f4..e68cdc27e1 100644 --- a/tests/out/glsl/boids.main.Compute.glsl +++ b/tests/out/glsl/boids.main.Compute.glsl @@ -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; diff --git a/tests/out/glsl/interpolate.frag_main.Fragment.glsl b/tests/out/glsl/interpolate.frag_main.Fragment.glsl index 918e78b9cf..1f35aebe6a 100644 --- a/tests/out/glsl/interpolate.frag_main.Fragment.glsl +++ b/tests/out/glsl/interpolate.frag_main.Fragment.glsl @@ -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; diff --git a/tests/out/glsl/interpolate.main.Fragment.glsl b/tests/out/glsl/interpolate.main.Fragment.glsl deleted file mode 100644 index 918e78b9cf..0000000000 --- a/tests/out/glsl/interpolate.main.Fragment.glsl +++ /dev/null @@ -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; -} - diff --git a/tests/out/glsl/interpolate.main.Vertex.glsl b/tests/out/glsl/interpolate.main.Vertex.glsl deleted file mode 100644 index 40d64ccf0d..0000000000 --- a/tests/out/glsl/interpolate.main.Vertex.glsl +++ /dev/null @@ -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; -} - diff --git a/tests/out/glsl/interpolate.vert_main.Vertex.glsl b/tests/out/glsl/interpolate.vert_main.Vertex.glsl index 40d64ccf0d..588b72ac64 100644 --- a/tests/out/glsl/interpolate.vert_main.Vertex.glsl +++ b/tests/out/glsl/interpolate.vert_main.Vertex.glsl @@ -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; diff --git a/tests/out/glsl/operators.main.Compute.glsl b/tests/out/glsl/operators.main.Compute.glsl index 39a284352b..9a9fa5c899 100644 --- a/tests/out/glsl/operators.main.Compute.glsl +++ b/tests/out/glsl/operators.main.Compute.glsl @@ -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)); diff --git a/tests/out/glsl/quad-vert.main.Vertex.glsl b/tests/out/glsl/quad-vert.main.Vertex.glsl index 26df417673..d3c740e05b 100644 --- a/tests/out/glsl/quad-vert.main.Vertex.glsl +++ b/tests/out/glsl/quad-vert.main.Vertex.glsl @@ -3,21 +3,21 @@ precision highp float; precision highp int; +struct gen_gl_PerVertex { + vec4 gen_gl_Position; + float gen_gl_PointSize; + float gen_gl_ClipDistance[1]; + float gen_gl_CullDistance[1]; +}; 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 { - vec4 gen_gl_Position; - float gen_gl_PointSize; - float gen_gl_ClipDistance[1]; - float gen_gl_CullDistance[1]; -} perVertexStruct; +gen_gl_PerVertex perVertexStruct; vec2 a_pos_1 = vec2(0.0, 0.0); diff --git a/tests/out/glsl/quad.frag_main.Fragment.glsl b/tests/out/glsl/quad.frag_main.Fragment.glsl index b290ed5a0e..637f890f75 100644 --- a/tests/out/glsl/quad.frag_main.Fragment.glsl +++ b/tests/out/glsl/quad.frag_main.Fragment.glsl @@ -7,7 +7,6 @@ struct VertexOutput { vec2 uv; vec4 position; }; - uniform highp sampler2D _group_0_binding_0; smooth in vec2 _vs2fs_location0; diff --git a/tests/out/glsl/quad.fs_extra.Fragment.glsl b/tests/out/glsl/quad.fs_extra.Fragment.glsl index 9770ffb60a..932fa76ba0 100644 --- a/tests/out/glsl/quad.fs_extra.Fragment.glsl +++ b/tests/out/glsl/quad.fs_extra.Fragment.glsl @@ -7,7 +7,6 @@ struct VertexOutput { vec2 uv; vec4 position; }; - layout(location = 0) out vec4 _fs2p_location0; void main() { diff --git a/tests/out/glsl/quad.vert_main.Vertex.glsl b/tests/out/glsl/quad.vert_main.Vertex.glsl index cda43f5086..5a1f432b28 100644 --- a/tests/out/glsl/quad.vert_main.Vertex.glsl +++ b/tests/out/glsl/quad.vert_main.Vertex.glsl @@ -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; diff --git a/tests/out/glsl/shadow.fs_main.Fragment.glsl b/tests/out/glsl/shadow.fs_main.Fragment.glsl index 7220591b92..594b250eae 100644 --- a/tests/out/glsl/shadow.fs_main.Fragment.glsl +++ b/tests/out/glsl/shadow.fs_main.Fragment.glsl @@ -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; diff --git a/tests/out/glsl/skybox.fs_main.Fragment.glsl b/tests/out/glsl/skybox.fs_main.Fragment.glsl index 08b733b848..916ea2c8ae 100644 --- a/tests/out/glsl/skybox.fs_main.Fragment.glsl +++ b/tests/out/glsl/skybox.fs_main.Fragment.glsl @@ -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; diff --git a/tests/out/glsl/skybox.vs_main.Vertex.glsl b/tests/out/glsl/skybox.vs_main.Vertex.glsl index 06494a70f5..f7490130a7 100644 --- a/tests/out/glsl/skybox.vs_main.Vertex.glsl +++ b/tests/out/glsl/skybox.vs_main.Vertex.glsl @@ -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; diff --git a/tests/out/ir/collatz.ron b/tests/out/ir/collatz.ron index bd2e15ee3f..5fff82ee3c 100644 --- a/tests/out/ir/collatz.ron +++ b/tests/out/ir/collatz.ron @@ -18,7 +18,6 @@ ( name: Some("PrimeIndices"), inner: Struct( - top_level: true, members: [ ( name: Some("data"), diff --git a/tests/out/ir/shadow.ron b/tests/out/ir/shadow.ron index 13220b7aea..0356f1371e 100644 --- a/tests/out/ir/shadow.ron +++ b/tests/out/ir/shadow.ron @@ -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"), diff --git a/tests/out/spv/access.spvasm b/tests/out/spv/access.spvasm index 96127247af..cd6fea5375 100644 --- a/tests/out/spv/access.spvasm +++ b/tests/out/spv/access.spvasm @@ -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 diff --git a/tests/out/spv/boids.spvasm b/tests/out/spv/boids.spvasm index 12d0cf69ae..c8f5894888 100644 --- a/tests/out/spv/boids.spvasm +++ b/tests/out/spv/boids.spvasm @@ -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 +%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 +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 -%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 +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 %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 +%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 -%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 +%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 \ No newline at end of file diff --git a/tests/out/spv/bounds-check-restrict.spvasm b/tests/out/spv/bounds-check-restrict.spvasm index 1f83ea976c..1927123467 100644 --- a/tests/out/spv/bounds-check-restrict.spvasm +++ b/tests/out/spv/bounds-check-restrict.spvasm @@ -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 diff --git a/tests/out/spv/bounds-check-zero.spvasm b/tests/out/spv/bounds-check-zero.spvasm index 82138feeb9..5e5aac07fa 100644 --- a/tests/out/spv/bounds-check-zero.spvasm +++ b/tests/out/spv/bounds-check-zero.spvasm @@ -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 diff --git a/tests/out/spv/collatz.spvasm b/tests/out/spv/collatz.spvasm index fd2065bb5a..bd2dd1f5fb 100644 --- a/tests/out/spv/collatz.spvasm +++ b/tests/out/spv/collatz.spvasm @@ -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 diff --git a/tests/out/spv/extra.spvasm b/tests/out/spv/extra.spvasm index 6f6f153a38..f17b286ac9 100644 --- a/tests/out/spv/extra.spvasm +++ b/tests/out/spv/extra.spvasm @@ -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 diff --git a/tests/out/spv/pointers.spvasm b/tests/out/spv/pointers.spvasm index 661a8e2c2e..968720f9d5 100644 --- a/tests/out/spv/pointers.spvasm +++ b/tests/out/spv/pointers.spvasm @@ -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 diff --git a/tests/out/spv/policy-mix.spvasm b/tests/out/spv/policy-mix.spvasm index 1e34e65e6f..685acfeaa2 100644 --- a/tests/out/spv/policy-mix.spvasm +++ b/tests/out/spv/policy-mix.spvasm @@ -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 \ No newline at end of file diff --git a/tests/out/spv/shadow.spvasm b/tests/out/spv/shadow.spvasm index a0554ed94a..3d4f4119a4 100644 --- a/tests/out/spv/shadow.spvasm +++ b/tests/out/spv/shadow.spvasm @@ -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 \ No newline at end of file diff --git a/tests/out/spv/skybox.spvasm b/tests/out/spv/skybox.spvasm index 16e597cb89..278ec0bb37 100644 --- a/tests/out/spv/skybox.spvasm +++ b/tests/out/spv/skybox.spvasm @@ -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 \ No newline at end of file diff --git a/tests/out/wgsl/210-bevy-2d-shader-frag.wgsl b/tests/out/wgsl/210-bevy-2d-shader-frag.wgsl index c578384dac..e66de740d6 100644 --- a/tests/out/wgsl/210-bevy-2d-shader-frag.wgsl +++ b/tests/out/wgsl/210-bevy-2d-shader-frag.wgsl @@ -1,4 +1,3 @@ -[[block]] struct ColorMaterial_color { Color: vec4; }; diff --git a/tests/out/wgsl/210-bevy-2d-shader-vert.wgsl b/tests/out/wgsl/210-bevy-2d-shader-vert.wgsl index 7db8c99255..c368e38d95 100644 --- a/tests/out/wgsl/210-bevy-2d-shader-vert.wgsl +++ b/tests/out/wgsl/210-bevy-2d-shader-vert.wgsl @@ -1,14 +1,11 @@ -[[block]] struct Camera { ViewProj: mat4x4; }; -[[block]] struct Transform { Model: mat4x4; }; -[[block]] struct Sprite_size { size: vec2; }; diff --git a/tests/out/wgsl/210-bevy-shader-vert.wgsl b/tests/out/wgsl/210-bevy-shader-vert.wgsl index 94ccaea3e9..4f476b9f15 100644 --- a/tests/out/wgsl/210-bevy-shader-vert.wgsl +++ b/tests/out/wgsl/210-bevy-shader-vert.wgsl @@ -1,9 +1,7 @@ -[[block]] struct Camera { ViewProj: mat4x4; }; -[[block]] struct Transform { Model: mat4x4; }; diff --git a/tests/out/wgsl/246-collatz-comp.wgsl b/tests/out/wgsl/246-collatz-comp.wgsl index e392695f7d..fc344ea509 100644 --- a/tests/out/wgsl/246-collatz-comp.wgsl +++ b/tests/out/wgsl/246-collatz-comp.wgsl @@ -1,4 +1,3 @@ -[[block]] struct PrimeIndices { indices: [[stride(4)]] array; }; diff --git a/tests/out/wgsl/800-out-of-bounds-panic-vert.wgsl b/tests/out/wgsl/800-out-of-bounds-panic-vert.wgsl index 06fd0f486d..351482d7da 100644 --- a/tests/out/wgsl/800-out-of-bounds-panic-vert.wgsl +++ b/tests/out/wgsl/800-out-of-bounds-panic-vert.wgsl @@ -1,9 +1,7 @@ -[[block]] struct Globals { view_matrix: mat4x4; }; -[[block]] struct VertexPushConstants { world_matrix: mat4x4; }; diff --git a/tests/out/wgsl/896-push-constant-vert.wgsl b/tests/out/wgsl/896-push-constant-vert.wgsl index 063f22a4f1..dec43d8d44 100644 --- a/tests/out/wgsl/896-push-constant-vert.wgsl +++ b/tests/out/wgsl/896-push-constant-vert.wgsl @@ -1,4 +1,3 @@ -[[block]] struct PushConstants { example: f32; }; diff --git a/tests/out/wgsl/access.wgsl b/tests/out/wgsl/access.wgsl index d56b4c5704..e437077157 100644 --- a/tests/out/wgsl/access.wgsl +++ b/tests/out/wgsl/access.wgsl @@ -1,4 +1,3 @@ -[[block]] struct Bar { matrix: mat4x4; atom: atomic; diff --git a/tests/out/wgsl/bevy-pbr-frag.wgsl b/tests/out/wgsl/bevy-pbr-frag.wgsl index 01b8120e62..f73f517d4e 100644 --- a/tests/out/wgsl/bevy-pbr-frag.wgsl +++ b/tests/out/wgsl/bevy-pbr-frag.wgsl @@ -9,17 +9,14 @@ struct DirectionalLight { color: vec4; }; -[[block]] struct CameraViewProj { ViewProj: mat4x4; }; -[[block]] struct CameraPosition { CameraPos: vec4; }; -[[block]] struct Lights { AmbientColor: vec4; NumLights: vec4; @@ -27,27 +24,22 @@ struct Lights { DirectionalLights: [[stride(32)]] array; }; -[[block]] struct StandardMaterial_base_color { base_color: vec4; }; -[[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; }; diff --git a/tests/out/wgsl/bevy-pbr-vert.wgsl b/tests/out/wgsl/bevy-pbr-vert.wgsl index 1863c9681b..34f836704e 100644 --- a/tests/out/wgsl/bevy-pbr-vert.wgsl +++ b/tests/out/wgsl/bevy-pbr-vert.wgsl @@ -1,9 +1,7 @@ -[[block]] struct CameraViewProj { ViewProj: mat4x4; }; -[[block]] struct Transform { Model: mat4x4; }; diff --git a/tests/out/wgsl/boids.wgsl b/tests/out/wgsl/boids.wgsl index 4dd4cd40ea..6a5d4c269b 100644 --- a/tests/out/wgsl/boids.wgsl +++ b/tests/out/wgsl/boids.wgsl @@ -3,7 +3,6 @@ struct Particle { vel: vec2; }; -[[block]] struct SimParams { deltaT: f32; rule1Distance: f32; @@ -14,7 +13,6 @@ struct SimParams { rule3Scale: f32; }; -[[block]] struct Particles { particles: [[stride(16)]] array; }; diff --git a/tests/out/wgsl/collatz.wgsl b/tests/out/wgsl/collatz.wgsl index a07b148fdf..281df6fb78 100644 --- a/tests/out/wgsl/collatz.wgsl +++ b/tests/out/wgsl/collatz.wgsl @@ -1,4 +1,3 @@ -[[block]] struct PrimeIndices { data: [[stride(4)]] array; }; diff --git a/tests/out/wgsl/constant-array-size-vert.wgsl b/tests/out/wgsl/constant-array-size-vert.wgsl index 268595b489..02d04dc520 100644 --- a/tests/out/wgsl/constant-array-size-vert.wgsl +++ b/tests/out/wgsl/constant-array-size-vert.wgsl @@ -1,4 +1,3 @@ -[[block]] struct Data { vecs: [[stride(16)]] array,42u>; }; diff --git a/tests/out/wgsl/empty-global-name.wgsl b/tests/out/wgsl/empty-global-name.wgsl index 1732aaf0b2..f912dfdff4 100644 --- a/tests/out/wgsl/empty-global-name.wgsl +++ b/tests/out/wgsl/empty-global-name.wgsl @@ -1,4 +1,3 @@ -[[block]] struct type_1 { member: i32; }; diff --git a/tests/out/wgsl/extra.wgsl b/tests/out/wgsl/extra.wgsl index 710519b8eb..b9fbfb74ca 100644 --- a/tests/out/wgsl/extra.wgsl +++ b/tests/out/wgsl/extra.wgsl @@ -1,4 +1,3 @@ -[[block]] struct PushConstants { index: u32; double: vec2; diff --git a/tests/out/wgsl/pointers.wgsl b/tests/out/wgsl/pointers.wgsl index c4c328fb4e..df8e6553ae 100644 --- a/tests/out/wgsl/pointers.wgsl +++ b/tests/out/wgsl/pointers.wgsl @@ -1,4 +1,3 @@ -[[block]] struct DynamicArray { arr: [[stride(4)]] array; }; diff --git a/tests/out/wgsl/quad-vert.wgsl b/tests/out/wgsl/quad-vert.wgsl index cff7f21775..b08aae9e68 100644 --- a/tests/out/wgsl/quad-vert.wgsl +++ b/tests/out/wgsl/quad-vert.wgsl @@ -1,4 +1,3 @@ -[[block]] struct gl_PerVertex { [[builtin(position)]] gl_Position: vec4; }; diff --git a/tests/out/wgsl/shadow.wgsl b/tests/out/wgsl/shadow.wgsl index c42d823e5b..d27b0a8951 100644 --- a/tests/out/wgsl/shadow.wgsl +++ b/tests/out/wgsl/shadow.wgsl @@ -1,4 +1,3 @@ -[[block]] struct Globals { num_lights: vec4; }; @@ -9,7 +8,6 @@ struct Light { color: vec4; }; -[[block]] struct Lights { data: [[stride(96)]] array; }; diff --git a/tests/out/wgsl/skybox.wgsl b/tests/out/wgsl/skybox.wgsl index 749725507b..8b52c97da0 100644 --- a/tests/out/wgsl/skybox.wgsl +++ b/tests/out/wgsl/skybox.wgsl @@ -3,7 +3,6 @@ struct VertexOutput { [[location(0)]] uv: vec3; }; -[[block]] struct Data { proj_inv: mat4x4; view: mat4x4; diff --git a/tests/wgsl-errors.rs b/tests/wgsl-errors.rs index 13b50109c0..7607359de3 100644 --- a/tests/wgsl-errors.rs +++ b/tests/wgsl-errors.rs @@ -804,17 +804,6 @@ fn invalid_arrays() { }) } - check_validation_error! { - r#" - [[block]] struct Block { value: f32; }; - type Bad = array; - "#: - Err(naga::valid::ValidationError::Type { - error: naga::valid::TypeError::NestedTopLevel, - .. - }) - } - check_validation_error! { r#" type Bad = [[stride(2)]] array; @@ -1154,7 +1143,6 @@ fn invalid_runtime_sized_arrays() { arr: array; }; - [[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; };