diff --git a/.gitignore b/.gitignore index 08a609ce9a..2aa145cd83 100644 --- a/.gitignore +++ b/.gitignore @@ -17,3 +17,4 @@ Cargo.lock /*.wgsl /*.hlsl /*.txt +*.code-workspace \ No newline at end of file diff --git a/src/back/msl/mod.rs b/src/back/msl/mod.rs index 819d48a8fd..1ca4584523 100644 --- a/src/back/msl/mod.rs +++ b/src/back/msl/mod.rs @@ -35,6 +35,8 @@ mod writer; pub use writer::Writer; +use self::writer::EntryPointInfo; + pub type Slot = u8; pub type InlineSamplerIndex = u8; @@ -79,6 +81,7 @@ pub struct EntryPointResources { pub type EntryPointResourceMap = std::collections::BTreeMap; +#[derive(Debug, Clone)] enum ResolvedBinding { BuiltIn(crate::BuiltIn), Attribute(u32), @@ -94,7 +97,7 @@ enum ResolvedBinding { Resource(BindTarget), } -#[derive(Copy, Clone)] +#[derive(Debug, Copy, Clone)] enum ResolvedInterpolation { CenterPerspective, CenterNoPerspective, @@ -494,7 +497,7 @@ pub struct TranslationInfo { /// corresponds to an entry point index. /// ///Note: Some entry points may fail translation because of missing bindings. - pub entry_point_names: Vec>, + pub entry_point_info: Vec>, } pub fn write_string( diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index 67ab887285..e9be2cd68e 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -1,10 +1,12 @@ -use super::{sampler as sm, Error, LocationMode, Options, PipelineOptions, TranslationInfo}; +use super::{ + sampler as sm, BindTarget, EntryPointError, Error, LocationMode, Options, PipelineOptions, + ResolvedBinding, TranslationInfo, +}; use crate::{ arena::Handle, - back, - proc::index, - proc::{self, NameKey, TypeResolution}, - valid, FastHashMap, FastHashSet, + back::{self, msl::BindSamplerTarget}, + proc::{self, index, NameKey, TypeResolution}, + valid, ArraySize, EntryPoint, FastHashMap, FastHashSet, GlobalVariable, Module, StorageAccess, }; use bit_set::BitSet; use std::{ @@ -82,8 +84,8 @@ struct TypeContext<'a> { handle: Handle, gctx: proc::GlobalCtx<'a>, names: &'a FastHashMap, - access: crate::StorageAccess, - binding: Option<&'a super::ResolvedBinding>, + access: StorageAccess, + binding: Option<&'a ResolvedBinding>, first_time: bool, } @@ -217,7 +219,7 @@ impl<'a> Display for TypeContext<'a> { ..*self }; - if let Some(&super::ResolvedBinding::Resource(super::BindTarget { + if let Some(&ResolvedBinding::Resource(BindTarget { binding_array_size: Some(override_size), .. })) = self.binding @@ -234,11 +236,11 @@ impl<'a> Display for TypeContext<'a> { } struct TypedGlobalVariable<'a> { - module: &'a crate::Module, + module: &'a Module, names: &'a FastHashMap, - handle: Handle, + handle: Handle, usage: valid::GlobalUse, - binding: Option<&'a super::ResolvedBinding>, + binding: Option<&'a ResolvedBinding>, reference: bool, } @@ -342,7 +344,7 @@ fn should_pack_struct_member( members: &[crate::StructMember], span: u32, index: usize, - module: &crate::Module, + module: &Module, ) -> Option { let member = &members[index]; //Note: this is imperfect - the same structure can be used for host-shared @@ -500,7 +502,7 @@ struct ExpressionContext<'a> { function: &'a crate::Function, origin: FunctionOrigin, info: &'a valid::FunctionInfo, - module: &'a crate::Module, + module: &'a Module, pipeline_options: &'a PipelineOptions, policies: index::BoundsCheckPolicies, @@ -1051,7 +1053,7 @@ impl Writer { /// dynamically sized array. fn put_dynamic_array_max_index( &mut self, - handle: Handle, + handle: Handle, context: &ExpressionContext, ) -> BackendResult { let global = &context.module.global_variables[handle]; @@ -1212,7 +1214,7 @@ impl Writer { fn put_const_expression( &mut self, expr_handle: Handle, - module: &crate::Module, + module: &Module, ) -> BackendResult { self.put_possibly_const_expression( expr_handle, @@ -1226,7 +1228,7 @@ impl Writer { &mut self, expr_handle: Handle, expressions: &crate::Arena, - module: &crate::Module, + module: &Module, put_expression: E, ) -> BackendResult where @@ -1417,8 +1419,19 @@ impl Writer { write!(self.out, "{name}")?; } crate::Expression::GlobalVariable(handle) => { + // If this variable belongs to a binding, and we are in the entry point, + // we must prefix it with the appropriate argument buffer name. + let variable = &context.module.global_variables[handle]; let name = &self.names[&NameKey::GlobalVariable(handle)]; - write!(self.out, "{name}")?; + + match (variable.binding.as_ref(), &context.origin) { + (Some(_), FunctionOrigin::EntryPoint(_)) => { + write!(self.out, "argumentBuffer.{name}")?; + } + _ => { + write!(self.out, "{name}")?; + } + } } crate::Expression::LocalVariable(handle) => { let name_key = match context.origin { @@ -2805,6 +2818,13 @@ impl Writer { } else { separate = true; } + + // If this global variable is part of a binding, prefix its name + // with the name of the corresponding argument buffer. + if let Some(_) = var.binding { + write!(self.out, "argumentBuffer.")?; + } + write!(self.out, "{name}")?; } supports_array_length |= @@ -3047,8 +3067,8 @@ impl Writer { pub fn write( &mut self, - module: &crate::Module, - info: &valid::ModuleInfo, + module: &Module, + module_info: &valid::ModuleInfo, options: &Options, pipeline_options: &PipelineOptions, ) -> Result { @@ -3078,28 +3098,7 @@ impl Writer { crate::TypeInner::RayQuery => true, _ => false, }) { - let tab = back::INDENT; - writeln!(self.out, "struct {RAY_QUERY_TYPE} {{")?; - let full_type = format!("{RT_NAMESPACE}::intersector<{RT_NAMESPACE}::instancing, {RT_NAMESPACE}::triangle_data, {RT_NAMESPACE}::world_space_data>"); - writeln!(self.out, "{tab}{full_type} {RAY_QUERY_FIELD_INTERSECTOR};")?; - writeln!( - self.out, - "{tab}{full_type}::result_type {RAY_QUERY_FIELD_INTERSECTION};" - )?; - writeln!(self.out, "{tab}bool {RAY_QUERY_FIELD_READY} = false;")?; - writeln!(self.out, "}};")?; - writeln!(self.out, "constexpr {NAMESPACE}::uint {RAY_QUERY_FUN_MAP_INTERSECTION}(const {RT_NAMESPACE}::intersection_type ty) {{")?; - let v_triangle = back::RayIntersectionType::Triangle as u32; - let v_bbox = back::RayIntersectionType::BoundingBox as u32; - writeln!( - self.out, - "{tab}return ty=={RT_NAMESPACE}::intersection_type::triangle ? {v_triangle} : " - )?; - writeln!( - self.out, - "{tab}{tab}ty=={RT_NAMESPACE}::intersection_type::bounding_box ? {v_bbox} : 0;" - )?; - writeln!(self.out, "}}")?; + self.write_ray_queries()?; } if options .bounds_check_policies @@ -3132,7 +3131,33 @@ impl Writer { self.write_type_defs(module)?; self.write_global_constants(module)?; - self.write_functions(module, info, options, pipeline_options) + self.write_functions(module, module_info, options, pipeline_options) + } + + fn write_ray_queries(&mut self) -> Result<(), Error> { + let tab = back::INDENT; + writeln!(self.out, "struct {RAY_QUERY_TYPE} {{")?; + let full_type = format!("{RT_NAMESPACE}::intersector<{RT_NAMESPACE}::instancing, {RT_NAMESPACE}::triangle_data, {RT_NAMESPACE}::world_space_data>"); + writeln!(self.out, "{tab}{full_type} {RAY_QUERY_FIELD_INTERSECTOR};")?; + writeln!( + self.out, + "{tab}{full_type}::result_type {RAY_QUERY_FIELD_INTERSECTION};" + )?; + writeln!(self.out, "{tab}bool {RAY_QUERY_FIELD_READY} = false;")?; + writeln!(self.out, "}};")?; + writeln!(self.out, "constexpr {NAMESPACE}::uint {RAY_QUERY_FUN_MAP_INTERSECTION}(const {RT_NAMESPACE}::intersection_type ty) {{")?; + let v_triangle = back::RayIntersectionType::Triangle as u32; + let v_bbox = back::RayIntersectionType::BoundingBox as u32; + writeln!( + self.out, + "{tab}return ty=={RT_NAMESPACE}::intersection_type::triangle ? {v_triangle} : " + )?; + writeln!( + self.out, + "{tab}{tab}ty=={RT_NAMESPACE}::intersection_type::bounding_box ? {v_bbox} : 0;" + )?; + writeln!(self.out, "}}")?; + Ok(()) } /// Write the definition for the `DefaultConstructible` class. @@ -3158,7 +3183,7 @@ impl Writer { Ok(()) } - fn write_type_defs(&mut self, module: &crate::Module) -> BackendResult { + fn write_type_defs(&mut self, module: &Module) -> BackendResult { for (handle, ty) in module.types.iter() { if !ty.needs_alias() { continue; @@ -3338,7 +3363,7 @@ impl Writer { } /// Writes all named constants - fn write_global_constants(&mut self, module: &crate::Module) -> BackendResult { + fn write_global_constants(&mut self, module: &Module) -> BackendResult { let constants = module.constants.iter().filter(|&(_, c)| c.name.is_some()); for (handle, constant) in constants { @@ -3440,12 +3465,14 @@ impl Writer { // Returns the array of mapped entry point names. fn write_functions( &mut self, - module: &crate::Module, - mod_info: &valid::ModuleInfo, + module: &Module, + module_info: &valid::ModuleInfo, options: &Options, pipeline_options: &PipelineOptions, ) -> Result { let mut pass_through_globals = Vec::new(); + + // Emit each user-defined function for (fun_handle, fun) in module.functions.iter() { log::trace!( "function {:?}, handle {:?}", @@ -3453,7 +3480,7 @@ impl Writer { fun_handle ); - let fun_info = &mod_info[fun_handle]; + let fun_info = &module_info[fun_handle]; pass_through_globals.clear(); let mut supports_array_length = false; for (handle, var) in module.global_variables.iter() { @@ -3571,7 +3598,7 @@ impl Writer { module, pipeline_options, }, - mod_info, + mod_info: module_info, result_struct: None, }; self.named_expressions.clear(); @@ -3580,293 +3607,318 @@ impl Writer { writeln!(self.out, "}}")?; } - let mut info = TranslationInfo { - entry_point_names: Vec::with_capacity(module.entry_points.len()), + let mut translation_info = TranslationInfo { + entry_point_info: Vec::with_capacity(module.entry_points.len()), }; + for (ep_index, ep) in module.entry_points.iter().enumerate() { - let fun = &ep.function; - let fun_info = mod_info.get_entry_point(ep_index); - let mut ep_error = None; + translation_info + .entry_point_info + .push(self.write_entry_point(CreateEntryPointInfo { + ep, + ep_index, + module, + mod_info: module_info, + options, + pipeline_options, + })?); + } + + Ok(translation_info) + } + + fn write_barrier(&mut self, flags: crate::Barrier, level: back::Level) -> BackendResult { + // Note: OR-ring bitflags requires `__HAVE_MEMFLAG_OPERATORS__`, + // so we try to avoid it here. + if flags.is_empty() { + writeln!( + self.out, + "{level}{NAMESPACE}::threadgroup_barrier({NAMESPACE}::mem_flags::mem_none);", + )?; + } + if flags.contains(crate::Barrier::STORAGE) { + writeln!( + self.out, + "{level}{NAMESPACE}::threadgroup_barrier({NAMESPACE}::mem_flags::mem_device);", + )?; + } + if flags.contains(crate::Barrier::WORK_GROUP) { + writeln!( + self.out, + "{level}{NAMESPACE}::threadgroup_barrier({NAMESPACE}::mem_flags::mem_threadgroup);", + )?; + } + Ok(()) + } + /// Write an argument buffer for an entry point. + /// + /// Iterates through the bindings used by entry point, generating + /// a struct whose members contain each binding for that group with a monotonically + /// increasing ID. + /// + /// For example, if we had some bind group: + /// + /// ```wgsl + /// @group(0) @binding(0) var u_texture : texture_2d; + /// @group(0) @binding(1) var u_sampler : sampler; + /// @group(1) @binding(0) var v_texture : texture_2d; + /// @group(1) @binding(0) var z_texture : texture_2d; + /// ``` + /// + /// This would generate: + /// + /// ```metal + /// struct ArgumentBufferGroup0 { + /// metal::texture2d u_texture [[id(0)]]; + /// metal::sampler u_sampler [[id(1)]]; + /// metal::texture2d v_texture [[id(2)]]; + /// metal::texture2d z_texture [[id(3)]]; + /// }; + /// ``` + /// + /// In the case of binding arrays, IDs will increase by the size of the array, as + /// required by Metal. + /// + /// For more information on Argument Buffers, see section 2.13 of + /// [the Metal Spec](https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf) + fn write_argument_buffer( + &mut self, + module: &Module, + mut members: Vec, + fun_name: &str, + ) -> Result { + // begin the argument buffer + writeln!(self.out, "struct {fun_name}ArgumentBuffer {{")?; + + log::trace!("Writing argument buffer for {fun_name}"); + + let mut argument_buffer_entries = ArgumentBufferEntries::default(); + + let mut id = 0; + for member in members.drain(..) { + let ArgumentBufferMember { + var, + handle, + resolved_binding, + } = member; + + // This must be valid. + let binding = var.binding.clone().unwrap(); + argument_buffer_entries.insert(binding, id); + + // write this member + let name = &self.names[&NameKey::GlobalVariable(handle)]; + + let gctx = module.to_ctx(); + let var_type = &gctx.types[var.ty]; log::trace!( - "entry point {:?}, index {:?}", - fun.name.as_deref().unwrap_or("(anonymous)"), - ep_index + "Writing argument buffer member {resolved_binding:?}: {name} - {var_type:?}" ); - // Is any global variable used by this entry point dynamically sized? - let supports_array_length = module - .global_variables - .iter() - .filter(|&(handle, _)| !fun_info[handle].is_empty()) - .any(|(_, var)| needs_array_length(var.ty, &module.types)); - - // skip this entry point if any global bindings are missing, - // or their types are incompatible. - if !options.fake_missing_bindings { - for (var_handle, var) in module.global_variables.iter() { - if fun_info[var_handle].is_empty() { - continue; - } - match var.space { - crate::AddressSpace::Uniform - | crate::AddressSpace::Storage { .. } - | crate::AddressSpace::Handle => { - let br = match var.binding { - Some(ref br) => br, - None => { - let var_name = var.name.clone().unwrap_or_default(); - ep_error = - Some(super::EntryPointError::MissingBinding(var_name)); - break; - } - }; - let target = options.get_resource_binding_target(ep, br); - let good = match target { - Some(target) => { - let binding_ty = match module.types[var.ty].inner { - crate::TypeInner::BindingArray { base, .. } => { - &module.types[base].inner - } - ref ty => ty, - }; - match *binding_ty { - crate::TypeInner::Image { .. } => target.texture.is_some(), - crate::TypeInner::Sampler { .. } => { - target.sampler.is_some() - } - _ => target.buffer.is_some(), - } - } - None => false, - }; - if !good { - ep_error = - Some(super::EntryPointError::MissingBindTarget(br.clone())); - break; - } - } - crate::AddressSpace::PushConstant => { - if let Err(e) = options.resolve_push_constants(ep) { - ep_error = Some(e); - break; - } - } - crate::AddressSpace::Function - | crate::AddressSpace::Private - | crate::AddressSpace::WorkGroup => {} - } - } - if supports_array_length { - if let Err(err) = options.resolve_sizes_buffer(ep) { - ep_error = Some(err); - } + // let id = + // get_argument_buffer_member_id(&resolved_binding, options.fake_missing_bindings)?; + + // Determine the storage access required for this variable + let storage_access = match var.space { + crate::AddressSpace::Storage { access } => access, + _ => match var_type.inner { + crate::TypeInner::Image { + class: crate::ImageClass::Storage { access, .. }, + .. + } => access, + crate::TypeInner::BindingArray { base, .. } => match gctx.types[base].inner { + crate::TypeInner::Image { + class: crate::ImageClass::Storage { access, .. }, + .. + } => access, + _ => crate::StorageAccess::LOAD, + }, + _ => crate::StorageAccess::LOAD, + }, + }; + + // Next, get the Metal name (eg. device, constant) for this variable's space + // and determine whether it needs to be qualified with "const". + let (space, access, reference) = match var.space.to_msl_name() { + Some(space) => { + let access = if var.space.needs_access_qualifier() + && !storage_access.contains(StorageAccess::STORE) + { + "const" + } else { + "" + }; + (space, access, "&") } - } + _ => ("", "", ""), + }; - if let Some(err) = ep_error { - info.entry_point_names.push(Err(err)); - continue; + // build the type name + let ty_name = TypeContext { + handle: var.ty, + gctx, + names: &self.names, + access: storage_access, + binding: Some(&resolved_binding), + first_time: false, + }; + + match writeln!( + self.out, + "{}{space}{}{ty_name}{}{access}{reference} {name} [[id({id})]];", + back::INDENT, + if space.is_empty() { "" } else { " " }, + if access.is_empty() { "" } else { " " }, + ) { + Err(e) => { + log::error!("Error writing argument buffer at member {var:?}"); + return Err(e.into()); + } + _ => {} } - let fun_name = &self.names[&NameKey::EntryPoint(ep_index as _)]; - info.entry_point_names.push(Ok(fun_name.clone())); - writeln!(self.out)?; + // If this binding was an array, we need to increment ID by the size of the array. - let (em_str, in_mode, out_mode) = match ep.stage { - crate::ShaderStage::Vertex => ( - "vertex", - LocationMode::VertexInput, - LocationMode::VertexOutput, - ), - crate::ShaderStage::Fragment { .. } => ( - "fragment", - LocationMode::FragmentInput, - LocationMode::FragmentOutput, - ), - crate::ShaderStage::Compute { .. } => { - ("kernel", LocationMode::Uniform, LocationMode::Uniform) + // Dynamically sized array bindings need to be treated as a special case: + match resolved_binding { + ResolvedBinding::Resource(BindTarget { + binding_array_size: Some(array_size), + .. + }) => { + id += array_size; + continue; } + _ => {} }; - // List all the Naga `EntryPoint`'s `Function`'s arguments, - // flattening structs into their members. In Metal, we will pass - // each of these values to the entry point as a separate argument— - // except for the varyings, handled next. - let mut flattened_arguments = Vec::new(); - for (arg_index, arg) in fun.arguments.iter().enumerate() { - match module.types[arg.ty].inner { - crate::TypeInner::Struct { ref members, .. } => { - for (member_index, member) in members.iter().enumerate() { - let member_index = member_index as u32; - flattened_arguments.push(( - NameKey::StructMember(arg.ty, member_index), - member.ty, - member.binding.as_ref(), - )); - } - } - _ => flattened_arguments.push(( - NameKey::EntryPointArgument(ep_index as _, arg_index as u32), - arg.ty, - arg.binding.as_ref(), - )), + // Otherwise, check the variable's type to see if it's an array. + match var_type.inner { + crate::TypeInner::Array { + size: ArraySize::Constant(size), + .. + } + | crate::TypeInner::BindingArray { + size: ArraySize::Constant(size), + .. + } => { + id += size.get(); } - } - // Identify the varyings among the argument values, and emit a - // struct type named `Input` to hold them. - let stage_in_name = format!("{fun_name}Input"); - let varyings_member_name = self.namer.call("varyings"); - let mut has_varyings = false; - if !flattened_arguments.is_empty() { - writeln!(self.out, "struct {stage_in_name} {{")?; - for &(ref name_key, ty, binding) in flattened_arguments.iter() { - let binding = match binding { - Some(ref binding @ &crate::Binding::Location { .. }) => binding, - _ => continue, - }; - has_varyings = true; - let name = &self.names[name_key]; - let ty_name = TypeContext { - handle: ty, - gctx: module.to_ctx(), - names: &self.names, - access: crate::StorageAccess::empty(), - binding: None, - first_time: false, - }; - let resolved = options.resolve_local_binding(binding, in_mode)?; - write!(self.out, "{}{} {}", back::INDENT, ty_name, name)?; - resolved.try_fmt(&mut self.out)?; - writeln!(self.out, ";")?; + // Not an array - just increment by 1. + _ => { + id += 1; } - writeln!(self.out, "}};")?; } + } - // Define a struct type named for the return value, if any, named - // `Output`. - let stage_out_name = format!("{fun_name}Output"); - let result_member_name = self.namer.call("member"); - let result_type_name = match fun.result { - Some(ref result) => { - let mut result_members = Vec::new(); - if let crate::TypeInner::Struct { ref members, .. } = - module.types[result.ty].inner - { - for (member_index, member) in members.iter().enumerate() { - result_members.push(( - &self.names[&NameKey::StructMember(result.ty, member_index as u32)], - member.ty, - member.binding.as_ref(), - )); - } - } else { - result_members.push(( - &result_member_name, - result.ty, - result.binding.as_ref(), - )); - } + // finish the struct + writeln!(self.out, "}};")?; - writeln!(self.out, "struct {stage_out_name} {{")?; - let mut has_point_size = false; - for (name, ty, binding) in result_members { - let ty_name = TypeContext { - handle: ty, - gctx: module.to_ctx(), - names: &self.names, - access: crate::StorageAccess::empty(), - binding: None, - first_time: true, - }; - let binding = binding.ok_or(Error::Validation)?; + Ok(argument_buffer_entries) + } - if let crate::Binding::BuiltIn(crate::BuiltIn::PointSize) = *binding { - has_point_size = true; - if !pipeline_options.allow_and_force_point_size { - continue; - } - } + /// Emits a given entry point by gathering up the statically accessed variables an functions + /// in that entry point + fn write_entry_point( + &mut self, + entry_point_info: CreateEntryPointInfo, + ) -> Result, Error> { + let CreateEntryPointInfo { + ep, + ep_index, + module, + mod_info, + options, + pipeline_options, + } = entry_point_info; + let fun = &ep.function; + let fun_info = mod_info.get_entry_point(ep_index); + + log::trace!( + "entry point {:?}, index {:?}", + fun.name.as_deref().unwrap_or("(anonymous)"), + ep_index + ); - let array_len = match module.types[ty].inner { - crate::TypeInner::Array { - size: crate::ArraySize::Constant(size), - .. - } => Some(size), - _ => None, - }; - let resolved = options.resolve_local_binding(binding, out_mode)?; - write!(self.out, "{}{} {}", back::INDENT, ty_name, name)?; - if let Some(array_len) = array_len { - write!(self.out, " [{array_len}]")?; - } - resolved.try_fmt(&mut self.out)?; - writeln!(self.out, ";")?; - } + // Is any global variable used by this entry point dynamically sized? + let has_dynamically_sized_global_variables = module + .global_variables + .iter() + .filter(|&(handle, _)| !fun_info[handle].is_empty()) + .any(|(_, var)| needs_array_length(var.ty, &module.types)); + + // If any global bindings are missing, or their types are incompatible, report this entry point + // as invalid. + let fun_name = self.names[&NameKey::EntryPoint(ep_index as _)].clone(); + if let Err(e) = validate_entry_point( + options, + module, + fun_info, + ep, + has_dynamically_sized_global_variables, + ) { + return Ok(Err(e)); + } - if pipeline_options.allow_and_force_point_size - && ep.stage == crate::ShaderStage::Vertex - && !has_point_size - { - // inject the point size output last - writeln!( - self.out, - "{}float _point_size [[point_size]];", - back::INDENT - )?; - } - writeln!(self.out, "}};")?; - &stage_out_name - } - None => "void", - }; + writeln!(self.out)?; - // Write the entry point function's name, and begin its argument list. - writeln!(self.out, "{em_str} {result_type_name} {fun_name}(")?; - let mut is_first_argument = true; + let (em_str, in_mode, out_mode) = match ep.stage { + crate::ShaderStage::Vertex => ( + "vertex", + LocationMode::VertexInput, + LocationMode::VertexOutput, + ), + crate::ShaderStage::Fragment { .. } => ( + "fragment", + LocationMode::FragmentInput, + LocationMode::FragmentOutput, + ), + crate::ShaderStage::Compute { .. } => { + ("kernel", LocationMode::Uniform, LocationMode::Uniform) + } + }; - // If we have produced a struct holding the `EntryPoint`'s - // `Function`'s arguments' varyings, pass that struct first. - if has_varyings { - writeln!( - self.out, - " {stage_in_name} {varyings_member_name} [[stage_in]]" - )?; - is_first_argument = false; + // List all the Naga `EntryPoint`'s `Function`'s arguments, + // flattening structs into their members. In Metal, we will pass + // each of these values to the entry point as a separate argument— + // except for the varyings, handled next. + // TODO: These could potentially be incorporated into the argument buffer. + let mut flattened_arguments = Vec::new(); + for (arg_index, arg) in fun.arguments.iter().enumerate() { + match module.types[arg.ty].inner { + crate::TypeInner::Struct { ref members, .. } => { + for (member_index, member) in members.iter().enumerate() { + let member_index = member_index as u32; + flattened_arguments.push(( + NameKey::StructMember(arg.ty, member_index), + member.ty, + member.binding.as_ref(), + )); + } + } + _ => flattened_arguments.push(( + NameKey::EntryPointArgument(ep_index as _, arg_index as u32), + arg.ty, + arg.binding.as_ref(), + )), } + } - let mut local_invocation_id = None; - - // Then pass the remaining arguments not included in the varyings - // struct. - // - // Since `Namer.reset` wasn't expecting struct members to be - // suddenly injected into the normal namespace like this, - // `self.names` doesn't keep them distinct from other variables. - // Generate fresh names for these arguments, and remember the - // mapping. - let mut flattened_member_names = FastHashMap::default(); + // Identify the varyings among the argument values, and emit a + // struct type named `Input` to hold them. + let stage_in_name = format!("{fun_name}Input"); + let varyings_member_name = self.namer.call("varyings"); + let mut has_varyings = false; + if !flattened_arguments.is_empty() { + writeln!(self.out, "struct {stage_in_name} {{")?; for &(ref name_key, ty, binding) in flattened_arguments.iter() { let binding = match binding { - Some(binding @ &crate::Binding::BuiltIn { .. }) => binding, + Some(ref binding @ &crate::Binding::Location { .. }) => binding, _ => continue, }; - let name = if let NameKey::StructMember(ty, index) = *name_key { - // We should always insert a fresh entry here, but use - // `or_insert` to get a reference to the `String` we just - // inserted. - flattened_member_names - .entry(NameKey::StructMember(ty, index)) - .or_insert_with(|| self.namer.call(&self.names[name_key])) - } else { - &self.names[name_key] - }; - - if binding == &crate::Binding::BuiltIn(crate::BuiltIn::LocalInvocationId) { - local_invocation_id = Some(name_key); - } - + has_varyings = true; + let name = &self.names[name_key]; let ty_name = TypeContext { handle: ty, gctx: module.to_ctx(), @@ -3876,308 +3928,608 @@ impl Writer { first_time: false, }; let resolved = options.resolve_local_binding(binding, in_mode)?; - let separator = if is_first_argument { - is_first_argument = false; - ' ' - } else { - ',' - }; - write!(self.out, "{separator} {ty_name} {name}")?; + write!(self.out, "{}{} {}", back::INDENT, ty_name, name)?; resolved.try_fmt(&mut self.out)?; - writeln!(self.out)?; + writeln!(self.out, ";")?; } + writeln!(self.out, "}};")?; + } - let need_workgroup_variables_initialization = - self.need_workgroup_variables_initialization(options, ep, module, fun_info); - - if need_workgroup_variables_initialization && local_invocation_id.is_none() { - let separator = if is_first_argument { - is_first_argument = false; - ' ' + // Define a struct type named for the return value, if any, named + // `Output`. + let stage_out_name = format!("{fun_name}Output"); + let result_member_name = self.namer.call("member"); + let result_type_name = match fun.result { + Some(ref result) => { + let mut result_members = Vec::new(); + if let crate::TypeInner::Struct { ref members, .. } = module.types[result.ty].inner + { + for (member_index, member) in members.iter().enumerate() { + result_members.push(( + &self.names[&NameKey::StructMember(result.ty, member_index as u32)], + member.ty, + member.binding.as_ref(), + )); + } } else { - ',' - }; - writeln!( - self.out, - "{separator} {NAMESPACE}::uint3 __local_invocation_id [[thread_position_in_threadgroup]]" - )?; - } - - // Those global variables used by this entry point and its callees - // get passed as arguments. `Private` globals are an exception, they - // don't outlive this invocation, so we declare them below as locals - // within the entry point. - for (handle, var) in module.global_variables.iter() { - let usage = fun_info[handle]; - if usage.is_empty() || var.space == crate::AddressSpace::Private { - continue; + result_members.push((&result_member_name, result.ty, result.binding.as_ref())); } - // the resolves have already been checked for `!fake_missing_bindings` case - let resolved = match var.space { - crate::AddressSpace::PushConstant => options.resolve_push_constants(ep).ok(), - crate::AddressSpace::WorkGroup => None, - crate::AddressSpace::Storage { .. } if options.lang_version < (2, 0) => { - return Err(Error::UnsupportedAddressSpace(var.space)) - } - _ => options - .resolve_resource_binding(ep, var.binding.as_ref().unwrap()) - .ok(), - }; - if let Some(ref resolved) = resolved { - // Inline samplers are be defined in the EP body - if resolved.as_inline_sampler(options).is_some() { - continue; + + writeln!(self.out, "struct {stage_out_name} {{")?; + let mut has_point_size = false; + for (name, ty, binding) in result_members { + let ty_name = TypeContext { + handle: ty, + gctx: module.to_ctx(), + names: &self.names, + access: crate::StorageAccess::empty(), + binding: None, + first_time: true, + }; + let binding = binding.ok_or(Error::Validation)?; + + if let crate::Binding::BuiltIn(crate::BuiltIn::PointSize) = *binding { + has_point_size = true; + if !pipeline_options.allow_and_force_point_size { + continue; + } } - } - let tyvar = TypedGlobalVariable { - module, - names: &self.names, - handle, - usage, - binding: resolved.as_ref(), - reference: true, - }; - let separator = if is_first_argument { - is_first_argument = false; - ' ' - } else { - ',' - }; - write!(self.out, "{separator} ")?; - tyvar.try_fmt(&mut self.out)?; - if let Some(resolved) = resolved { + let array_len = match module.types[ty].inner { + crate::TypeInner::Array { + size: crate::ArraySize::Constant(size), + .. + } => Some(size), + _ => None, + }; + let resolved = options.resolve_local_binding(binding, out_mode)?; + write!(self.out, "{}{} {}", back::INDENT, ty_name, name)?; + if let Some(array_len) = array_len { + write!(self.out, " [{array_len}]")?; + } resolved.try_fmt(&mut self.out)?; + writeln!(self.out, ";")?; } - if let Some(value) = var.init { - write!(self.out, " = ")?; - self.put_const_expression(value, module)?; + + if pipeline_options.allow_and_force_point_size + && ep.stage == crate::ShaderStage::Vertex + && !has_point_size + { + // inject the point size output last + writeln!( + self.out, + "{}float _point_size [[point_size]];", + back::INDENT + )?; } - writeln!(self.out)?; + writeln!(self.out, "}};")?; + &stage_out_name } + None => "void", + }; - // If this entry uses any variable-length arrays, their sizes are - // passed as a final struct-typed argument. - if supports_array_length { - // this is checked earlier - let resolved = options.resolve_sizes_buffer(ep).unwrap(); - let separator = if module.global_variables.is_empty() { - ' ' - } else { - ',' - }; - write!( - self.out, - "{separator} constant _mslBufferSizes& _buffer_sizes", - )?; - resolved.try_fmt(&mut self.out)?; - writeln!(self.out)?; + let mut global_variables_passed_by_argument = Vec::new(); + let mut argument_buffer_members = Vec::new(); + + // Walk through each of the global variables used within this entry point and + // determine how they should be handled. + // + // - Resource bindings are collected into an argument buffer + // - Private global variables and inline samplers are declared as locals within the entry point + // - All other global variables are passed as arguments to the entry point + for (handle, var) in module.global_variables.iter() { + let usage = fun_info[handle]; + if usage.is_empty() || var.space == crate::AddressSpace::Private { + continue; } + // the resolves have already been checked for `!fake_missing_bindings` case + let resolved = match var.space { + crate::AddressSpace::PushConstant => options.resolve_push_constants(ep).ok(), + crate::AddressSpace::WorkGroup => None, + crate::AddressSpace::Storage { .. } if options.lang_version < (2, 0) => { + return Err(Error::UnsupportedAddressSpace(var.space)) + } + _ => { + // This global variable refers to a resource binding, which will require access through + // an argument buffer. - // end of the entry point argument list - writeln!(self.out, ") {{")?; + // TODO: Check for case of writable texture on tier 1 target. - if need_workgroup_variables_initialization { - self.write_workgroup_variables_initialization( - module, - mod_info, - fun_info, - local_invocation_id, - )?; - } + // These unwraps are safe as all bindings been validated when beginning to write this entry point. + let resolved_binding = options + .resolve_resource_binding(ep, var.binding.as_ref().unwrap()) + .unwrap(); + argument_buffer_members.push(ArgumentBufferMember { + var, + handle, + resolved_binding, + }); - // Metal doesn't support private mutable variables outside of functions, - // so we put them here, just like the locals. - for (handle, var) in module.global_variables.iter() { - let usage = fun_info[handle]; - if usage.is_empty() { continue; } - if var.space == crate::AddressSpace::Private { - let tyvar = TypedGlobalVariable { - module, - names: &self.names, - handle, - usage, - binding: None, - reference: false, - }; - write!(self.out, "{}", back::INDENT)?; - tyvar.try_fmt(&mut self.out)?; - match var.init { - Some(value) => { - write!(self.out, " = ")?; - self.put_const_expression(value, module)?; - writeln!(self.out, ";")?; - } - None => { - writeln!(self.out, " = {{}};")?; - } - }; - } else if let Some(ref binding) = var.binding { - // write an inline sampler - let resolved = options.resolve_resource_binding(ep, binding).unwrap(); - if let Some(sampler) = resolved.as_inline_sampler(options) { - let name = &self.names[&NameKey::GlobalVariable(handle)]; - writeln!( - self.out, - "{}constexpr {}::sampler {}(", - back::INDENT, - NAMESPACE, - name - )?; - self.put_inline_sampler_properties(back::Level(2), sampler)?; - writeln!(self.out, "{});", back::INDENT)?; - } + }; + + if let Some(ref resolved) = resolved { + // Inline samplers are defined in the body of the entry point + if resolved.as_inline_sampler(options).is_some() { + continue; } } - // Now take the arguments that we gathered into structs, and the - // structs that we flattened into arguments, and emit local - // variables with initializers that put everything back the way the - // body code expects. - // - // If we had to generate fresh names for struct members passed as - // arguments, be sure to use those names when rebuilding the struct. - // - // "Each day, I change some zeros to ones, and some ones to zeros. - // The rest, I leave alone." - for (arg_index, arg) in fun.arguments.iter().enumerate() { - let arg_name = - &self.names[&NameKey::EntryPointArgument(ep_index as _, arg_index as u32)]; - match module.types[arg.ty].inner { - crate::TypeInner::Struct { ref members, .. } => { - let struct_name = &self.names[&NameKey::Type(arg.ty)]; - write!( - self.out, - "{}const {} {} = {{ ", - back::INDENT, - struct_name, - arg_name - )?; - for (member_index, member) in members.iter().enumerate() { - let key = NameKey::StructMember(arg.ty, member_index as u32); - // If it's not in the varying struct, then we should - // have passed it as its own argument and assigned - // it a new name. - let name = match member.binding { - Some(crate::Binding::BuiltIn { .. }) => { - &flattened_member_names[&key] - } - _ => &self.names[&key], - }; - if member_index != 0 { - write!(self.out, ", ")?; - } - if let Some(crate::Binding::Location { .. }) = member.binding { - write!(self.out, "{varyings_member_name}.")?; - } - write!(self.out, "{name}")?; - } - writeln!(self.out, " }};")?; - } - _ => { - if let Some(crate::Binding::Location { .. }) = arg.binding { - writeln!( - self.out, - "{}const auto {} = {}.{};", - back::INDENT, - arg_name, - varyings_member_name, - arg_name - )?; - } - } - } + global_variables_passed_by_argument.push(GlobalVariablePassedByArgument { + var, + handle, + resolved_binding: resolved, + usage, + }); + } + + // If we have bindings that can be represented in an argument buffer, emit a struct + // that contains those bindings. + let argument_buffer_entries = + self.write_argument_buffer(module, argument_buffer_members, &fun_name)?; + + // Write the entry point function's name, and begin its argument list. + writeln!(self.out, "{em_str} {result_type_name} {fun_name}(")?; + let mut is_first_argument = true; + + // If we have produced a struct holding the `EntryPoint`'s + // `Function`'s arguments' varyings, pass that struct first. + if has_varyings { + writeln!( + self.out, + " {stage_in_name} {varyings_member_name} [[stage_in]]" + )?; + is_first_argument = false; + } + + let mut local_invocation_id = None; + + // Then pass the remaining arguments not included in the varyings + // struct. + // + // Since `Namer.reset` wasn't expecting struct members to be + // suddenly injected into the normal namespace like this, + // `self.names` doesn't keep them distinct from other variables. + // Generate fresh names for these arguments, and remember the + // mapping. + let mut flattened_member_names = FastHashMap::default(); + for &(ref name_key, ty, binding) in flattened_arguments.iter() { + let binding = match binding { + Some(binding @ &crate::Binding::BuiltIn { .. }) => binding, + _ => continue, + }; + let name = if let NameKey::StructMember(ty, index) = *name_key { + // We should always insert a fresh entry here, but use + // `or_insert` to get a reference to the `String` we just + // inserted. + flattened_member_names + .entry(NameKey::StructMember(ty, index)) + .or_insert_with(|| self.namer.call(&self.names[name_key])) + } else { + &self.names[name_key] + }; + + if binding == &crate::Binding::BuiltIn(crate::BuiltIn::LocalInvocationId) { + local_invocation_id = Some(name_key); } - // Finally, declare all the local variables that we need - //TODO: we can postpone this till the relevant expressions are emitted - for (local_handle, local) in fun.local_variables.iter() { - let name = &self.names[&NameKey::EntryPointLocal(ep_index as _, local_handle)]; - let ty_name = TypeContext { - handle: local.ty, - gctx: module.to_ctx(), + let ty_name = TypeContext { + handle: ty, + gctx: module.to_ctx(), + names: &self.names, + access: crate::StorageAccess::empty(), + binding: None, + first_time: false, + }; + let resolved = options.resolve_local_binding(binding, in_mode)?; + let separator = if is_first_argument { + is_first_argument = false; + ' ' + } else { + ',' + }; + write!(self.out, "{separator} {ty_name} {name}")?; + resolved.try_fmt(&mut self.out)?; + writeln!(self.out)?; + } + + let need_workgroup_variables_initialization = + self.need_workgroup_variables_initialization(options, ep, module, fun_info); + + if need_workgroup_variables_initialization && local_invocation_id.is_none() { + let separator = if is_first_argument { + is_first_argument = false; + ' ' + } else { + ',' + }; + writeln!( + self.out, + "{separator} {NAMESPACE}::uint3 __local_invocation_id [[thread_position_in_threadgroup]]" + )?; + } + + // If we have collected bindings into an argument buffer, pass that buffer as an argument to + // the entry-point. + if !argument_buffer_entries.is_empty() { + let separator = if is_first_argument { + is_first_argument = false; + ' ' + } else { + ',' + }; + write!(self.out, "{separator} ")?; + // TODO: Determine if this is a vertex entry point; if so, this should be passed in as const + writeln!( + self.out, + "device {fun_name}ArgumentBuffer& argumentBuffer [[buffer(0)]]" + )?; + } + + // All other global variables must be passed by argument + for global_variable in global_variables_passed_by_argument.drain(..) { + let GlobalVariablePassedByArgument { + var, + handle, + resolved_binding, + usage, + } = global_variable; + + let tyvar = TypedGlobalVariable { + module, + names: &self.names, + handle, + usage, + binding: resolved_binding.as_ref(), + reference: true, + }; + + let separator = if is_first_argument { + is_first_argument = false; + ' ' + } else { + ',' + }; + write!(self.out, "{separator} ")?; + tyvar.try_fmt(&mut self.out)?; + if let Some(resolved) = resolved_binding { + resolved.try_fmt(&mut self.out)?; + } + if let Some(value) = var.init { + write!(self.out, " = ")?; + self.put_const_expression(value, module)?; + } + writeln!(self.out)?; + } + + // If this entry uses any variable-length arrays, their sizes are + // passed as a final struct-typed argument. + if has_dynamically_sized_global_variables { + // this is checked earlier + let resolved = options.resolve_sizes_buffer(ep).unwrap(); + let separator = if module.global_variables.is_empty() { + ' ' + } else { + ',' + }; + write!( + self.out, + "{separator} constant _mslBufferSizes& _buffer_sizes", + )?; + resolved.try_fmt(&mut self.out)?; + writeln!(self.out)?; + } + + // end of the entry point argument list + writeln!(self.out, ") {{")?; + + if need_workgroup_variables_initialization { + self.write_workgroup_variables_initialization( + module, + mod_info, + fun_info, + local_invocation_id, + )?; + } + + // Metal doesn't support private mutable variables outside of functions, + // so we put them here, just like the locals. + for (handle, var) in module.global_variables.iter() { + let usage = fun_info[handle]; + if usage.is_empty() { + continue; + } + if var.space == crate::AddressSpace::Private { + let tyvar = TypedGlobalVariable { + module, names: &self.names, - access: crate::StorageAccess::empty(), + handle, + usage, binding: None, - first_time: false, + reference: false, }; - write!(self.out, "{}{} {}", back::INDENT, ty_name, name)?; - match local.init { + write!(self.out, "{}", back::INDENT)?; + tyvar.try_fmt(&mut self.out)?; + match var.init { Some(value) => { write!(self.out, " = ")?; self.put_const_expression(value, module)?; + writeln!(self.out, ";")?; } None => { - write!(self.out, " = {{}}")?; + writeln!(self.out, " = {{}};")?; } }; - writeln!(self.out, ";")?; + } else if let Some(ref binding) = var.binding { + // write an inline sampler + let resolved = options.resolve_resource_binding(ep, binding).unwrap(); + if let Some(sampler) = resolved.as_inline_sampler(options) { + let name = &self.names[&NameKey::GlobalVariable(handle)]; + writeln!( + self.out, + "{}constexpr {}::sampler {}(", + back::INDENT, + NAMESPACE, + name + )?; + self.put_inline_sampler_properties(back::Level(2), sampler)?; + writeln!(self.out, "{});", back::INDENT)?; + } } + } - let guarded_indices = - index::find_checked_indexes(module, fun, fun_info, options.bounds_check_policies); + // Now take the arguments that we gathered into structs, and the + // structs that we flattened into arguments, and emit local + // variables with initializers that put everything back the way the + // body code expects. + // + // If we had to generate fresh names for struct members passed as + // arguments, be sure to use those names when rebuilding the struct. + // + // "Each day, I change some zeros to ones, and some ones to zeros. + // The rest, I leave alone." + // + // TODO: This could potentially be simplified through argument buffers. + for (arg_index, arg) in fun.arguments.iter().enumerate() { + let arg_name = + &self.names[&NameKey::EntryPointArgument(ep_index as _, arg_index as u32)]; + match module.types[arg.ty].inner { + crate::TypeInner::Struct { ref members, .. } => { + let struct_name = &self.names[&NameKey::Type(arg.ty)]; + write!( + self.out, + "{}const {} {} = {{ ", + back::INDENT, + struct_name, + arg_name + )?; + for (member_index, member) in members.iter().enumerate() { + let key = NameKey::StructMember(arg.ty, member_index as u32); + // If it's not in the varying struct, then we should + // have passed it as its own argument and assigned + // it a new name. + let name = match member.binding { + Some(crate::Binding::BuiltIn { .. }) => &flattened_member_names[&key], + _ => &self.names[&key], + }; + if member_index != 0 { + write!(self.out, ", ")?; + } + if let Some(crate::Binding::Location { .. }) = member.binding { + write!(self.out, "{varyings_member_name}.")?; + } + write!(self.out, "{name}")?; + } + writeln!(self.out, " }};")?; + } + _ => { + if let Some(crate::Binding::Location { .. }) = arg.binding { + writeln!( + self.out, + "{}const auto {} = {}.{};", + back::INDENT, + arg_name, + varyings_member_name, + arg_name + )?; + } + } + } + } - let context = StatementContext { - expression: ExpressionContext { - function: fun, - origin: FunctionOrigin::EntryPoint(ep_index as _), - info: fun_info, - policies: options.bounds_check_policies, - guarded_indices, - module, - pipeline_options, - }, - mod_info, - result_struct: Some(&stage_out_name), + // Finally, declare all the local variables that we need + //TODO: we can postpone this till the relevant expressions are emitted + for (local_handle, local) in fun.local_variables.iter() { + let name = &self.names[&NameKey::EntryPointLocal(ep_index as _, local_handle)]; + let ty_name = TypeContext { + handle: local.ty, + gctx: module.to_ctx(), + names: &self.names, + access: crate::StorageAccess::empty(), + binding: None, + first_time: false, }; - self.named_expressions.clear(); - self.update_expressions_to_bake(fun, fun_info, &context.expression); - self.put_block(back::Level(1), &fun.body, &context)?; - writeln!(self.out, "}}")?; - if ep_index + 1 != module.entry_points.len() { - writeln!(self.out)?; - } + write!(self.out, "{}{} {}", back::INDENT, ty_name, name)?; + match local.init { + Some(value) => { + write!(self.out, " = ")?; + self.put_const_expression(value, module)?; + } + None => { + write!(self.out, " = {{}}")?; + } + }; + writeln!(self.out, ";")?; } - Ok(info) + let guarded_indices = + index::find_checked_indexes(module, fun, fun_info, options.bounds_check_policies); + + let context = StatementContext { + expression: ExpressionContext { + function: fun, + origin: FunctionOrigin::EntryPoint(ep_index as _), + info: fun_info, + policies: options.bounds_check_policies, + guarded_indices, + module, + pipeline_options, + }, + mod_info, + result_struct: Some(&stage_out_name), + }; + self.named_expressions.clear(); + self.update_expressions_to_bake(fun, fun_info, &context.expression); + self.put_block(back::Level(1), &fun.body, &context)?; + writeln!(self.out, "}}")?; + if ep_index + 1 != module.entry_points.len() { + writeln!(self.out)?; + } + + Ok(Ok(EntryPointInfo { + name: fun_name, + argument_buffer_entries, + })) } +} - fn write_barrier(&mut self, flags: crate::Barrier, level: back::Level) -> BackendResult { - // Note: OR-ring bitflags requires `__HAVE_MEMFLAG_OPERATORS__`, - // so we try to avoid it here. - if flags.is_empty() { - writeln!( - self.out, - "{level}{NAMESPACE}::threadgroup_barrier({NAMESPACE}::mem_flags::mem_none);", - )?; - } - if flags.contains(crate::Barrier::STORAGE) { - writeln!( - self.out, - "{level}{NAMESPACE}::threadgroup_barrier({NAMESPACE}::mem_flags::mem_device);", - )?; +fn get_argument_buffer_member_id( + resolved_binding: &ResolvedBinding, + fake_missing_bindings: bool, +) -> Result { + let bind_target = match resolved_binding { + ResolvedBinding::Resource(b) => b, + // If we are faking missing bindings, then simply use the fake index we were given. + ResolvedBinding::User { index, .. } if fake_missing_bindings => return Ok(*index as _), + // It should not be possible for us to have been passed a variable that isn't resolved, + // so throw a validation error here. + _ => return Err(Error::Validation), + }; + + match bind_target { + BindTarget { + buffer: Some(slot), .. + } => Ok(*slot), + BindTarget { + texture: Some(slot), + .. + } => Ok(*slot), + BindTarget { + sampler: Some(BindSamplerTarget::Resource(slot)), + .. + } => Ok(*slot), + // It should not be possible for us to have been passed a variable without a valid bind + // target so throw a validation error here. + _ => Err(Error::Validation), + } +} + +struct CreateEntryPointInfo<'a> { + ep: &'a EntryPoint, + ep_index: usize, + module: &'a Module, + mod_info: &'a valid::ModuleInfo, + options: &'a Options, + pipeline_options: &'a PipelineOptions, +} + + +#[derive(Debug, Clone)] +pub struct EntryPointInfo { + pub name: String, + pub argument_buffer_entries: ArgumentBufferEntries, +} + +struct ArgumentBufferMember<'a> { + var: &'a GlobalVariable, + handle: Handle, + resolved_binding: ResolvedBinding, +} + +struct GlobalVariablePassedByArgument<'a> { + var: &'a GlobalVariable, + handle: Handle, + resolved_binding: Option, + usage: valid::GlobalUse, +} + +type ArgumentBufferEntries = FastHashMap; + +/// Check this entry point to determine if any global bindings are missing, or their types are +/// incompatible +fn validate_entry_point( + options: &Options, + module: &Module, + fun_info: &valid::FunctionInfo, + ep: &EntryPoint, + supports_array_length: bool, +) -> Result<(), EntryPointError> { + // If we're faking missing bindings, then there's no need to validate them. + if options.fake_missing_bindings { + return Ok(()); + } + + for (var_handle, var) in module.global_variables.iter() { + if fun_info[var_handle].is_empty() { + continue; } - if flags.contains(crate::Barrier::WORK_GROUP) { - writeln!( - self.out, - "{level}{NAMESPACE}::threadgroup_barrier({NAMESPACE}::mem_flags::mem_threadgroup);", - )?; + match var.space { + crate::AddressSpace::Uniform + | crate::AddressSpace::Storage { .. } + | crate::AddressSpace::Handle => { + let br = match var.binding { + Some(ref br) => br, + None => { + let var_name = var.name.clone().unwrap_or_default(); + return Err(EntryPointError::MissingBinding(var_name)); + } + }; + let target = options.get_resource_binding_target(ep, br); + let good = match target { + Some(target) => { + let binding_ty = match module.types[var.ty].inner { + crate::TypeInner::BindingArray { base, .. } => { + &module.types[base].inner + } + ref ty => ty, + }; + match *binding_ty { + crate::TypeInner::Image { .. } => target.texture.is_some(), + crate::TypeInner::Sampler { .. } => target.sampler.is_some(), + _ => target.buffer.is_some(), + } + } + None => false, + }; + if !good { + return Err(EntryPointError::MissingBindTarget(br.clone())); + } + } + crate::AddressSpace::PushConstant => { + options.resolve_push_constants(ep)?; + } + crate::AddressSpace::Function + | crate::AddressSpace::Private + | crate::AddressSpace::WorkGroup => {} } - Ok(()) } + if supports_array_length { + options.resolve_sizes_buffer(ep)?; + } + + Ok(()) } /// Initializing workgroup variables is more tricky for Metal because we have to deal /// with atomics at the type-level (which don't have a copy constructor). mod workgroup_mem_init { - use crate::EntryPoint; + use crate::{EntryPoint, GlobalVariable, Module}; use super::*; enum Access { - GlobalVariable(Handle), + GlobalVariable(Handle), StructMember(Handle, u32), Array(usize), } @@ -4247,7 +4599,7 @@ mod workgroup_mem_init { &mut self, options: &Options, ep: &EntryPoint, - module: &crate::Module, + module: &Module, fun_info: &valid::FunctionInfo, ) -> bool { options.zero_initialize_workgroup_memory @@ -4259,7 +4611,7 @@ mod workgroup_mem_init { pub(super) fn write_workgroup_variables_initialization( &mut self, - module: &crate::Module, + module: &Module, module_info: &valid::ModuleInfo, fun_info: &valid::FunctionInfo, local_invocation_id: Option<&NameKey>, @@ -4301,7 +4653,7 @@ mod workgroup_mem_init { fn write_workgroup_variable_initialization( &mut self, - module: &crate::Module, + module: &Module, module_info: &valid::ModuleInfo, ty: Handle, access_stack: &mut AccessStack, @@ -4373,7 +4725,7 @@ mod workgroup_mem_init { fn test_stack_size() { use crate::valid::{Capabilities, ValidationFlags}; // create a module with at least one expression nested - let mut module = crate::Module::default(); + let mut module = Module::default(); let mut fun = crate::Function::default(); let const_expr = fun.expressions.append( crate::Expression::Literal(crate::Literal::F32(1.0)), diff --git a/tests/out/msl/access.msl b/tests/out/msl/access.msl index e5d875dd19..e66102a435 100644 --- a/tests/out/msl/access.msl +++ b/tests/out/msl/access.msl @@ -168,12 +168,15 @@ struct foo_vertInput { struct foo_vertOutput { metal::float4 member [[position]]; }; +struct foo_vertArgumentBuffer { + device Bar& bar [[id(0)]]; + constant Baz& baz [[id(1)]]; + device metal::int2& qux [[id(2)]]; + constant MatCx2InArray& nested_mat_cx2_ [[id(3)]]; +}; vertex foo_vertOutput foo_vert( uint vi [[vertex_id]] -, device Bar const& bar [[buffer(0)]] -, constant Baz& baz [[buffer(1)]] -, device metal::int2 const& qux [[buffer(2)]] -, constant MatCx2InArray& nested_mat_cx2_ [[buffer(3)]] +, device foo_vertArgumentBuffer& argumentBuffer [[buffer(0)]] , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]] ) { float foo = {}; @@ -181,13 +184,13 @@ vertex foo_vertOutput foo_vert( foo = 0.0; float baz_1 = foo; foo = 1.0; - test_matrix_within_struct_accesses(baz); - test_matrix_within_array_within_struct_accesses(nested_mat_cx2_); - metal::float4x3 _matrix = bar._matrix; - type_9 arr_1 = bar.arr; - float b = bar._matrix[3].x; - int a_1 = bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value; - metal::int2 c = qux; + test_matrix_within_struct_accesses(argumentBuffer.baz); + test_matrix_within_array_within_struct_accesses(argumentBuffer.nested_mat_cx2_); + metal::float4x3 _matrix = argumentBuffer.bar._matrix; + type_9 arr_1 = argumentBuffer.bar.arr; + float b = argumentBuffer.bar._matrix[3].x; + int a_1 = argumentBuffer.bar.data[(1 + (_buffer_sizes.size1 - 160 - 8) / 8) - 2u].value; + metal::int2 c = argumentBuffer.qux; float _e34 = read_from_private(foo); c2_ = type_22 {a_1, static_cast(b), 3, 4, 5}; c2_.inner[vi + 1u] = 42; @@ -200,16 +203,19 @@ vertex foo_vertOutput foo_vert( struct foo_fragOutput { metal::float4 member_1 [[color(0)]]; }; +struct foo_fragArgumentBuffer { + device Bar& bar [[id(0)]]; + device metal::int2& qux [[id(1)]]; +}; fragment foo_fragOutput foo_frag( - device Bar& bar [[buffer(0)]] -, device metal::int2& qux [[buffer(2)]] + device foo_fragArgumentBuffer& argumentBuffer [[buffer(0)]] , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]] ) { - bar._matrix[1].z = 1.0; - bar._matrix = metal::float4x3(metal::float3(0.0), metal::float3(1.0), metal::float3(2.0), metal::float3(3.0)); - bar.arr = type_9 {metal::uint2(0u), metal::uint2(1u)}; - bar.data[1].value = 1; - qux = metal::int2 {}; + argumentBuffer.bar._matrix[1].z = 1.0; + argumentBuffer.bar._matrix = metal::float4x3(metal::float3(0.0), metal::float3(1.0), metal::float3(2.0), metal::float3(3.0)); + argumentBuffer.bar.arr = type_9 {metal::uint2(0u), metal::uint2(1u)}; + argumentBuffer.bar.data[1].value = 1; + argumentBuffer.qux = metal::int2 {}; return foo_fragOutput { metal::float4(0.0) }; } diff --git a/tests/out/msl/array-in-ctor.msl b/tests/out/msl/array-in-ctor.msl index 9428cb1e74..98b1a97ffe 100644 --- a/tests/out/msl/array-in-ctor.msl +++ b/tests/out/msl/array-in-ctor.msl @@ -11,8 +11,11 @@ struct Ah { type_1 inner; }; +struct cs_mainArgumentBuffer { + device Ah const& ah [[id(0)]]; +}; kernel void cs_main( - device Ah const& ah [[user(fake0)]] + device cs_mainArgumentBuffer& argumentBuffer [[buffer(0)]] ) { - Ah ah_1 = ah; + Ah ah_1 = argumentBuffer.ah; } diff --git a/tests/out/msl/atomicOps.msl b/tests/out/msl/atomicOps.msl index b7264e883d..32bf698eda 100644 --- a/tests/out/msl/atomicOps.msl +++ b/tests/out/msl/atomicOps.msl @@ -14,11 +14,14 @@ struct Struct { struct cs_mainInput { }; +struct cs_mainArgumentBuffer { + device metal::atomic_uint& storage_atomic_scalar [[id(0)]]; + device type_2& storage_atomic_arr [[id(1)]]; + device Struct& storage_struct [[id(3)]]; +}; kernel void cs_main( metal::uint3 id [[thread_position_in_threadgroup]] -, device metal::atomic_uint& storage_atomic_scalar [[user(fake0)]] -, device type_2& storage_atomic_arr [[user(fake0)]] -, device Struct& storage_struct [[user(fake0)]] +, device cs_mainArgumentBuffer& argumentBuffer [[buffer(0)]] , threadgroup metal::atomic_uint& workgroup_atomic_scalar , threadgroup type_2& workgroup_atomic_arr , threadgroup Struct& workgroup_struct @@ -34,90 +37,90 @@ kernel void cs_main( } } metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - metal::atomic_store_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); - metal::atomic_store_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); - metal::atomic_store_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); - metal::atomic_store_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + metal::atomic_store_explicit(&argumentBuffer.storage_atomic_scalar, 1u, metal::memory_order_relaxed); + metal::atomic_store_explicit(&argumentBuffer.storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + metal::atomic_store_explicit(&argumentBuffer.storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + metal::atomic_store_explicit(&argumentBuffer.storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); metal::atomic_store_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); metal::atomic_store_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); metal::atomic_store_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); metal::atomic_store_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - uint l0_ = metal::atomic_load_explicit(&storage_atomic_scalar, metal::memory_order_relaxed); - int l1_ = metal::atomic_load_explicit(&storage_atomic_arr.inner[1], metal::memory_order_relaxed); - uint l2_ = metal::atomic_load_explicit(&storage_struct.atomic_scalar, metal::memory_order_relaxed); - int l3_ = metal::atomic_load_explicit(&storage_struct.atomic_arr.inner[1], metal::memory_order_relaxed); + uint l0_ = metal::atomic_load_explicit(&argumentBuffer.storage_atomic_scalar, metal::memory_order_relaxed); + int l1_ = metal::atomic_load_explicit(&argumentBuffer.storage_atomic_arr.inner[1], metal::memory_order_relaxed); + uint l2_ = metal::atomic_load_explicit(&argumentBuffer.storage_struct.atomic_scalar, metal::memory_order_relaxed); + int l3_ = metal::atomic_load_explicit(&argumentBuffer.storage_struct.atomic_arr.inner[1], metal::memory_order_relaxed); uint l4_ = metal::atomic_load_explicit(&workgroup_atomic_scalar, metal::memory_order_relaxed); int l5_ = metal::atomic_load_explicit(&workgroup_atomic_arr.inner[1], metal::memory_order_relaxed); uint l6_ = metal::atomic_load_explicit(&workgroup_struct.atomic_scalar, metal::memory_order_relaxed); int l7_ = metal::atomic_load_explicit(&workgroup_struct.atomic_arr.inner[1], metal::memory_order_relaxed); metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - uint _e59 = metal::atomic_fetch_add_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); - int _e64 = metal::atomic_fetch_add_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); - uint _e68 = metal::atomic_fetch_add_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); - int _e74 = metal::atomic_fetch_add_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e59 = metal::atomic_fetch_add_explicit(&argumentBuffer.storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e64 = metal::atomic_fetch_add_explicit(&argumentBuffer.storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e68 = metal::atomic_fetch_add_explicit(&argumentBuffer.storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e74 = metal::atomic_fetch_add_explicit(&argumentBuffer.storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e77 = metal::atomic_fetch_add_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); int _e82 = metal::atomic_fetch_add_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e86 = metal::atomic_fetch_add_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); int _e92 = metal::atomic_fetch_add_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - uint _e95 = metal::atomic_fetch_sub_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); - int _e100 = metal::atomic_fetch_sub_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); - uint _e104 = metal::atomic_fetch_sub_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); - int _e110 = metal::atomic_fetch_sub_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e95 = metal::atomic_fetch_sub_explicit(&argumentBuffer.storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e100 = metal::atomic_fetch_sub_explicit(&argumentBuffer.storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e104 = metal::atomic_fetch_sub_explicit(&argumentBuffer.storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e110 = metal::atomic_fetch_sub_explicit(&argumentBuffer.storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e113 = metal::atomic_fetch_sub_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); int _e118 = metal::atomic_fetch_sub_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e122 = metal::atomic_fetch_sub_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); int _e128 = metal::atomic_fetch_sub_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - uint _e131 = metal::atomic_fetch_max_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); - int _e136 = metal::atomic_fetch_max_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); - uint _e140 = metal::atomic_fetch_max_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); - int _e146 = metal::atomic_fetch_max_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e131 = metal::atomic_fetch_max_explicit(&argumentBuffer.storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e136 = metal::atomic_fetch_max_explicit(&argumentBuffer.storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e140 = metal::atomic_fetch_max_explicit(&argumentBuffer.storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e146 = metal::atomic_fetch_max_explicit(&argumentBuffer.storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e149 = metal::atomic_fetch_max_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); int _e154 = metal::atomic_fetch_max_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e158 = metal::atomic_fetch_max_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); int _e164 = metal::atomic_fetch_max_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - uint _e167 = metal::atomic_fetch_min_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); - int _e172 = metal::atomic_fetch_min_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); - uint _e176 = metal::atomic_fetch_min_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); - int _e182 = metal::atomic_fetch_min_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e167 = metal::atomic_fetch_min_explicit(&argumentBuffer.storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e172 = metal::atomic_fetch_min_explicit(&argumentBuffer.storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e176 = metal::atomic_fetch_min_explicit(&argumentBuffer.storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e182 = metal::atomic_fetch_min_explicit(&argumentBuffer.storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e185 = metal::atomic_fetch_min_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); int _e190 = metal::atomic_fetch_min_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e194 = metal::atomic_fetch_min_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); int _e200 = metal::atomic_fetch_min_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - uint _e203 = metal::atomic_fetch_and_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); - int _e208 = metal::atomic_fetch_and_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); - uint _e212 = metal::atomic_fetch_and_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); - int _e218 = metal::atomic_fetch_and_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e203 = metal::atomic_fetch_and_explicit(&argumentBuffer.storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e208 = metal::atomic_fetch_and_explicit(&argumentBuffer.storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e212 = metal::atomic_fetch_and_explicit(&argumentBuffer.storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e218 = metal::atomic_fetch_and_explicit(&argumentBuffer.storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e221 = metal::atomic_fetch_and_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); int _e226 = metal::atomic_fetch_and_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e230 = metal::atomic_fetch_and_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); int _e236 = metal::atomic_fetch_and_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - uint _e239 = metal::atomic_fetch_or_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); - int _e244 = metal::atomic_fetch_or_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); - uint _e248 = metal::atomic_fetch_or_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); - int _e254 = metal::atomic_fetch_or_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e239 = metal::atomic_fetch_or_explicit(&argumentBuffer.storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e244 = metal::atomic_fetch_or_explicit(&argumentBuffer.storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e248 = metal::atomic_fetch_or_explicit(&argumentBuffer.storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e254 = metal::atomic_fetch_or_explicit(&argumentBuffer.storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e257 = metal::atomic_fetch_or_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); int _e262 = metal::atomic_fetch_or_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e266 = metal::atomic_fetch_or_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); int _e272 = metal::atomic_fetch_or_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); - uint _e275 = metal::atomic_fetch_xor_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); - int _e280 = metal::atomic_fetch_xor_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); - uint _e284 = metal::atomic_fetch_xor_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); - int _e290 = metal::atomic_fetch_xor_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e275 = metal::atomic_fetch_xor_explicit(&argumentBuffer.storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e280 = metal::atomic_fetch_xor_explicit(&argumentBuffer.storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e284 = metal::atomic_fetch_xor_explicit(&argumentBuffer.storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e290 = metal::atomic_fetch_xor_explicit(&argumentBuffer.storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e293 = metal::atomic_fetch_xor_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); int _e298 = metal::atomic_fetch_xor_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e302 = metal::atomic_fetch_xor_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); int _e308 = metal::atomic_fetch_xor_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); - uint _e311 = metal::atomic_exchange_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed); - int _e316 = metal::atomic_exchange_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); - uint _e320 = metal::atomic_exchange_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); - int _e326 = metal::atomic_exchange_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e311 = metal::atomic_exchange_explicit(&argumentBuffer.storage_atomic_scalar, 1u, metal::memory_order_relaxed); + int _e316 = metal::atomic_exchange_explicit(&argumentBuffer.storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed); + uint _e320 = metal::atomic_exchange_explicit(&argumentBuffer.storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed); + int _e326 = metal::atomic_exchange_explicit(&argumentBuffer.storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e329 = metal::atomic_exchange_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed); int _e334 = metal::atomic_exchange_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed); uint _e338 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed); diff --git a/tests/out/msl/binding-arrays.msl b/tests/out/msl/binding-arrays.msl index 10a5397910..2dbea5d86c 100644 --- a/tests/out/msl/binding-arrays.msl +++ b/tests/out/msl/binding-arrays.msl @@ -23,24 +23,27 @@ struct main_Input { struct main_Output { metal::float4 member [[color(0)]]; }; +struct main_ArgumentBuffer { + metal::array, 10> texture_array_unbounded [[id(0)]]; + metal::array, 5> texture_array_bounded [[id(10)]]; + metal::array, 5> texture_array_2darray [[id(15)]]; + metal::array, 5> texture_array_multisampled [[id(20)]]; + metal::array, 5> texture_array_depth [[id(25)]]; + metal::array, 5> texture_array_storage [[id(30)]]; + metal::array samp [[id(35)]]; + metal::array samp_comp [[id(40)]]; + constant UniformIndex& uni [[id(45)]]; +}; fragment main_Output main_( main_Input varyings [[stage_in]] -, metal::array, 10> texture_array_unbounded [[texture(0)]] -, metal::array, 5> texture_array_bounded [[user(fake0)]] -, metal::array, 5> texture_array_2darray [[user(fake0)]] -, metal::array, 5> texture_array_multisampled [[user(fake0)]] -, metal::array, 5> texture_array_depth [[user(fake0)]] -, metal::array, 5> texture_array_storage [[user(fake0)]] -, metal::array samp [[user(fake0)]] -, metal::array samp_comp [[user(fake0)]] -, constant UniformIndex& uni [[user(fake0)]] +, device main_ArgumentBuffer& argumentBuffer [[buffer(0)]] ) { const FragmentIn fragment_in = { varyings.index }; uint u1_ = {}; metal::uint2 u2_ = {}; float v1_ = {}; metal::float4 v4_ = {}; - uint uniform_index = uni.index; + uint uniform_index = argumentBuffer.uni.index; uint non_uniform_index = fragment_in.index; u1_ = 0u; u2_ = metal::uint2(0u); @@ -49,121 +52,121 @@ fragment main_Output main_( metal::float2 uv = metal::float2(0.0); metal::int2 pix = metal::int2(0); metal::uint2 _e23 = u2_; - u2_ = _e23 + metal::uint2(texture_array_unbounded[0].get_width(), texture_array_unbounded[0].get_height()); + u2_ = _e23 + metal::uint2(argumentBuffer.texture_array_unbounded[0].get_width(), argumentBuffer.texture_array_unbounded[0].get_height()); metal::uint2 _e28 = u2_; - u2_ = _e28 + metal::uint2(texture_array_unbounded[uniform_index].get_width(), texture_array_unbounded[uniform_index].get_height()); + u2_ = _e28 + metal::uint2(argumentBuffer.texture_array_unbounded[uniform_index].get_width(), argumentBuffer.texture_array_unbounded[uniform_index].get_height()); metal::uint2 _e33 = u2_; - u2_ = _e33 + metal::uint2(texture_array_unbounded[non_uniform_index].get_width(), texture_array_unbounded[non_uniform_index].get_height()); - metal::float4 _e42 = texture_array_bounded[0].gather(samp[0], uv); + u2_ = _e33 + metal::uint2(argumentBuffer.texture_array_unbounded[non_uniform_index].get_width(), argumentBuffer.texture_array_unbounded[non_uniform_index].get_height()); + metal::float4 _e42 = argumentBuffer.texture_array_bounded[0].gather(argumentBuffer.samp[0], uv); metal::float4 _e43 = v4_; v4_ = _e43 + _e42; - metal::float4 _e50 = texture_array_bounded[uniform_index].gather(samp[uniform_index], uv); + metal::float4 _e50 = argumentBuffer.texture_array_bounded[uniform_index].gather(argumentBuffer.samp[uniform_index], uv); metal::float4 _e51 = v4_; v4_ = _e51 + _e50; - metal::float4 _e58 = texture_array_bounded[non_uniform_index].gather(samp[non_uniform_index], uv); + metal::float4 _e58 = argumentBuffer.texture_array_bounded[non_uniform_index].gather(argumentBuffer.samp[non_uniform_index], uv); metal::float4 _e59 = v4_; v4_ = _e59 + _e58; - metal::float4 _e68 = texture_array_depth[0].gather_compare(samp_comp[0], uv, 0.0); + metal::float4 _e68 = argumentBuffer.texture_array_depth[0].gather_compare(argumentBuffer.samp_comp[0], uv, 0.0); metal::float4 _e69 = v4_; v4_ = _e69 + _e68; - metal::float4 _e76 = texture_array_depth[uniform_index].gather_compare(samp_comp[uniform_index], uv, 0.0); + metal::float4 _e76 = argumentBuffer.texture_array_depth[uniform_index].gather_compare(argumentBuffer.samp_comp[uniform_index], uv, 0.0); metal::float4 _e77 = v4_; v4_ = _e77 + _e76; - metal::float4 _e84 = texture_array_depth[non_uniform_index].gather_compare(samp_comp[non_uniform_index], uv, 0.0); + metal::float4 _e84 = argumentBuffer.texture_array_depth[non_uniform_index].gather_compare(argumentBuffer.samp_comp[non_uniform_index], uv, 0.0); metal::float4 _e85 = v4_; v4_ = _e85 + _e84; - metal::float4 _e91 = (uint(0) < texture_array_unbounded[0].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(texture_array_unbounded[0].get_width(0), texture_array_unbounded[0].get_height(0))) ? texture_array_unbounded[0].read(metal::uint2(pix), 0): DefaultConstructible()); + metal::float4 _e91 = (uint(0) < argumentBuffer.texture_array_unbounded[0].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(argumentBuffer.texture_array_unbounded[0].get_width(0), argumentBuffer.texture_array_unbounded[0].get_height(0))) ? argumentBuffer.texture_array_unbounded[0].read(metal::uint2(pix), 0): DefaultConstructible()); metal::float4 _e92 = v4_; v4_ = _e92 + _e91; - metal::float4 _e97 = (uint(0) < texture_array_unbounded[uniform_index].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(texture_array_unbounded[uniform_index].get_width(0), texture_array_unbounded[uniform_index].get_height(0))) ? texture_array_unbounded[uniform_index].read(metal::uint2(pix), 0): DefaultConstructible()); + metal::float4 _e97 = (uint(0) < argumentBuffer.texture_array_unbounded[uniform_index].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(argumentBuffer.texture_array_unbounded[uniform_index].get_width(0), argumentBuffer.texture_array_unbounded[uniform_index].get_height(0))) ? argumentBuffer.texture_array_unbounded[uniform_index].read(metal::uint2(pix), 0): DefaultConstructible()); metal::float4 _e98 = v4_; v4_ = _e98 + _e97; - metal::float4 _e103 = (uint(0) < texture_array_unbounded[non_uniform_index].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(texture_array_unbounded[non_uniform_index].get_width(0), texture_array_unbounded[non_uniform_index].get_height(0))) ? texture_array_unbounded[non_uniform_index].read(metal::uint2(pix), 0): DefaultConstructible()); + metal::float4 _e103 = (uint(0) < argumentBuffer.texture_array_unbounded[non_uniform_index].get_num_mip_levels() && metal::all(metal::uint2(pix) < metal::uint2(argumentBuffer.texture_array_unbounded[non_uniform_index].get_width(0), argumentBuffer.texture_array_unbounded[non_uniform_index].get_height(0))) ? argumentBuffer.texture_array_unbounded[non_uniform_index].read(metal::uint2(pix), 0): DefaultConstructible()); metal::float4 _e104 = v4_; v4_ = _e104 + _e103; uint _e110 = u1_; - u1_ = _e110 + texture_array_2darray[0].get_array_size(); + u1_ = _e110 + argumentBuffer.texture_array_2darray[0].get_array_size(); uint _e115 = u1_; - u1_ = _e115 + texture_array_2darray[uniform_index].get_array_size(); + u1_ = _e115 + argumentBuffer.texture_array_2darray[uniform_index].get_array_size(); uint _e120 = u1_; - u1_ = _e120 + texture_array_2darray[non_uniform_index].get_array_size(); + u1_ = _e120 + argumentBuffer.texture_array_2darray[non_uniform_index].get_array_size(); uint _e126 = u1_; - u1_ = _e126 + texture_array_bounded[0].get_num_mip_levels(); + u1_ = _e126 + argumentBuffer.texture_array_bounded[0].get_num_mip_levels(); uint _e131 = u1_; - u1_ = _e131 + texture_array_bounded[uniform_index].get_num_mip_levels(); + u1_ = _e131 + argumentBuffer.texture_array_bounded[uniform_index].get_num_mip_levels(); uint _e136 = u1_; - u1_ = _e136 + texture_array_bounded[non_uniform_index].get_num_mip_levels(); + u1_ = _e136 + argumentBuffer.texture_array_bounded[non_uniform_index].get_num_mip_levels(); uint _e142 = u1_; - u1_ = _e142 + texture_array_multisampled[0].get_num_samples(); + u1_ = _e142 + argumentBuffer.texture_array_multisampled[0].get_num_samples(); uint _e147 = u1_; - u1_ = _e147 + texture_array_multisampled[uniform_index].get_num_samples(); + u1_ = _e147 + argumentBuffer.texture_array_multisampled[uniform_index].get_num_samples(); uint _e152 = u1_; - u1_ = _e152 + texture_array_multisampled[non_uniform_index].get_num_samples(); - metal::float4 _e160 = texture_array_bounded[0].sample(samp[0], uv); + u1_ = _e152 + argumentBuffer.texture_array_multisampled[non_uniform_index].get_num_samples(); + metal::float4 _e160 = argumentBuffer.texture_array_bounded[0].sample(argumentBuffer.samp[0], uv); metal::float4 _e161 = v4_; v4_ = _e161 + _e160; - metal::float4 _e167 = texture_array_bounded[uniform_index].sample(samp[uniform_index], uv); + metal::float4 _e167 = argumentBuffer.texture_array_bounded[uniform_index].sample(argumentBuffer.samp[uniform_index], uv); metal::float4 _e168 = v4_; v4_ = _e168 + _e167; - metal::float4 _e174 = texture_array_bounded[non_uniform_index].sample(samp[non_uniform_index], uv); + metal::float4 _e174 = argumentBuffer.texture_array_bounded[non_uniform_index].sample(argumentBuffer.samp[non_uniform_index], uv); metal::float4 _e175 = v4_; v4_ = _e175 + _e174; - metal::float4 _e184 = texture_array_bounded[0].sample(samp[0], uv, metal::bias(0.0)); + metal::float4 _e184 = argumentBuffer.texture_array_bounded[0].sample(argumentBuffer.samp[0], uv, metal::bias(0.0)); metal::float4 _e185 = v4_; v4_ = _e185 + _e184; - metal::float4 _e192 = texture_array_bounded[uniform_index].sample(samp[uniform_index], uv, metal::bias(0.0)); + metal::float4 _e192 = argumentBuffer.texture_array_bounded[uniform_index].sample(argumentBuffer.samp[uniform_index], uv, metal::bias(0.0)); metal::float4 _e193 = v4_; v4_ = _e193 + _e192; - metal::float4 _e200 = texture_array_bounded[non_uniform_index].sample(samp[non_uniform_index], uv, metal::bias(0.0)); + metal::float4 _e200 = argumentBuffer.texture_array_bounded[non_uniform_index].sample(argumentBuffer.samp[non_uniform_index], uv, metal::bias(0.0)); metal::float4 _e201 = v4_; v4_ = _e201 + _e200; - float _e210 = texture_array_depth[0].sample_compare(samp_comp[0], uv, 0.0); + float _e210 = argumentBuffer.texture_array_depth[0].sample_compare(argumentBuffer.samp_comp[0], uv, 0.0); float _e211 = v1_; v1_ = _e211 + _e210; - float _e218 = texture_array_depth[uniform_index].sample_compare(samp_comp[uniform_index], uv, 0.0); + float _e218 = argumentBuffer.texture_array_depth[uniform_index].sample_compare(argumentBuffer.samp_comp[uniform_index], uv, 0.0); float _e219 = v1_; v1_ = _e219 + _e218; - float _e226 = texture_array_depth[non_uniform_index].sample_compare(samp_comp[non_uniform_index], uv, 0.0); + float _e226 = argumentBuffer.texture_array_depth[non_uniform_index].sample_compare(argumentBuffer.samp_comp[non_uniform_index], uv, 0.0); float _e227 = v1_; v1_ = _e227 + _e226; - float _e236 = texture_array_depth[0].sample_compare(samp_comp[0], uv, 0.0); + float _e236 = argumentBuffer.texture_array_depth[0].sample_compare(argumentBuffer.samp_comp[0], uv, 0.0); float _e237 = v1_; v1_ = _e237 + _e236; - float _e244 = texture_array_depth[uniform_index].sample_compare(samp_comp[uniform_index], uv, 0.0); + float _e244 = argumentBuffer.texture_array_depth[uniform_index].sample_compare(argumentBuffer.samp_comp[uniform_index], uv, 0.0); float _e245 = v1_; v1_ = _e245 + _e244; - float _e252 = texture_array_depth[non_uniform_index].sample_compare(samp_comp[non_uniform_index], uv, 0.0); + float _e252 = argumentBuffer.texture_array_depth[non_uniform_index].sample_compare(argumentBuffer.samp_comp[non_uniform_index], uv, 0.0); float _e253 = v1_; v1_ = _e253 + _e252; - metal::float4 _e261 = texture_array_bounded[0].sample(samp[0], uv, metal::gradient2d(uv, uv)); + metal::float4 _e261 = argumentBuffer.texture_array_bounded[0].sample(argumentBuffer.samp[0], uv, metal::gradient2d(uv, uv)); metal::float4 _e262 = v4_; v4_ = _e262 + _e261; - metal::float4 _e268 = texture_array_bounded[uniform_index].sample(samp[uniform_index], uv, metal::gradient2d(uv, uv)); + metal::float4 _e268 = argumentBuffer.texture_array_bounded[uniform_index].sample(argumentBuffer.samp[uniform_index], uv, metal::gradient2d(uv, uv)); metal::float4 _e269 = v4_; v4_ = _e269 + _e268; - metal::float4 _e275 = texture_array_bounded[non_uniform_index].sample(samp[non_uniform_index], uv, metal::gradient2d(uv, uv)); + metal::float4 _e275 = argumentBuffer.texture_array_bounded[non_uniform_index].sample(argumentBuffer.samp[non_uniform_index], uv, metal::gradient2d(uv, uv)); metal::float4 _e276 = v4_; v4_ = _e276 + _e275; - metal::float4 _e285 = texture_array_bounded[0].sample(samp[0], uv, metal::level(0.0)); + metal::float4 _e285 = argumentBuffer.texture_array_bounded[0].sample(argumentBuffer.samp[0], uv, metal::level(0.0)); metal::float4 _e286 = v4_; v4_ = _e286 + _e285; - metal::float4 _e293 = texture_array_bounded[uniform_index].sample(samp[uniform_index], uv, metal::level(0.0)); + metal::float4 _e293 = argumentBuffer.texture_array_bounded[uniform_index].sample(argumentBuffer.samp[uniform_index], uv, metal::level(0.0)); metal::float4 _e294 = v4_; v4_ = _e294 + _e293; - metal::float4 _e301 = texture_array_bounded[non_uniform_index].sample(samp[non_uniform_index], uv, metal::level(0.0)); + metal::float4 _e301 = argumentBuffer.texture_array_bounded[non_uniform_index].sample(argumentBuffer.samp[non_uniform_index], uv, metal::level(0.0)); metal::float4 _e302 = v4_; v4_ = _e302 + _e301; metal::float4 _e307 = v4_; - if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[0].get_width(), texture_array_storage[0].get_height()))) { - texture_array_storage[0].write(_e307, metal::uint2(pix)); + if (metal::all(metal::uint2(pix) < metal::uint2(argumentBuffer.texture_array_storage[0].get_width(), argumentBuffer.texture_array_storage[0].get_height()))) { + argumentBuffer.texture_array_storage[0].write(_e307, metal::uint2(pix)); } metal::float4 _e310 = v4_; - if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[uniform_index].get_width(), texture_array_storage[uniform_index].get_height()))) { - texture_array_storage[uniform_index].write(_e310, metal::uint2(pix)); + if (metal::all(metal::uint2(pix) < metal::uint2(argumentBuffer.texture_array_storage[uniform_index].get_width(), argumentBuffer.texture_array_storage[uniform_index].get_height()))) { + argumentBuffer.texture_array_storage[uniform_index].write(_e310, metal::uint2(pix)); } metal::float4 _e313 = v4_; - if (metal::all(metal::uint2(pix) < metal::uint2(texture_array_storage[non_uniform_index].get_width(), texture_array_storage[non_uniform_index].get_height()))) { - texture_array_storage[non_uniform_index].write(_e313, metal::uint2(pix)); + if (metal::all(metal::uint2(pix) < metal::uint2(argumentBuffer.texture_array_storage[non_uniform_index].get_width(), argumentBuffer.texture_array_storage[non_uniform_index].get_height()))) { + argumentBuffer.texture_array_storage[non_uniform_index].write(_e313, metal::uint2(pix)); } metal::uint2 _e314 = u2_; uint _e315 = u1_; diff --git a/tests/out/msl/boids.msl b/tests/out/msl/boids.msl index 1a81aaf684..600d874293 100644 --- a/tests/out/msl/boids.msl +++ b/tests/out/msl/boids.msl @@ -30,11 +30,14 @@ constant uint NUM_PARTICLES = 1500u; struct main_Input { }; +struct main_ArgumentBuffer { + constant SimParams& params [[id(0)]]; + device Particles const& particlesSrc [[id(1)]]; + device Particles& particlesDst [[id(2)]]; +}; kernel void main_( metal::uint3 global_invocation_id [[thread_position_in_grid]] -, constant SimParams& params [[buffer(0)]] -, device Particles const& particlesSrc [[buffer(1)]] -, device Particles& particlesDst [[buffer(2)]] +, device main_ArgumentBuffer& argumentBuffer [[buffer(0)]] , constant _mslBufferSizes& _buffer_sizes [[buffer(3)]] ) { metal::float2 vPos = {}; @@ -51,9 +54,9 @@ kernel void main_( if (index >= NUM_PARTICLES) { return; } - metal::float2 _e8 = particlesSrc.particles[index].pos; + metal::float2 _e8 = argumentBuffer.particlesSrc.particles[index].pos; vPos = _e8; - metal::float2 _e14 = particlesSrc.particles[index].vel; + metal::float2 _e14 = argumentBuffer.particlesSrc.particles[index].vel; vVel = _e14; cMass = metal::float2(0.0, 0.0); cVel = metal::float2(0.0, 0.0); @@ -77,14 +80,14 @@ kernel void main_( continue; } uint _e43 = i; - metal::float2 _e46 = particlesSrc.particles[_e43].pos; + metal::float2 _e46 = argumentBuffer.particlesSrc.particles[_e43].pos; pos = _e46; uint _e49 = i; - metal::float2 _e52 = particlesSrc.particles[_e49].vel; + metal::float2 _e52 = argumentBuffer.particlesSrc.particles[_e49].vel; vel = _e52; metal::float2 _e53 = pos; metal::float2 _e54 = vPos; - float _e58 = params.rule1Distance; + float _e58 = argumentBuffer.params.rule1Distance; if (metal::distance(_e53, _e54) < _e58) { metal::float2 _e60 = cMass; metal::float2 _e61 = pos; @@ -94,7 +97,7 @@ kernel void main_( } metal::float2 _e66 = pos; metal::float2 _e67 = vPos; - float _e71 = params.rule2Distance; + float _e71 = argumentBuffer.params.rule2Distance; if (metal::distance(_e66, _e67) < _e71) { metal::float2 _e73 = colVel; metal::float2 _e74 = pos; @@ -103,7 +106,7 @@ kernel void main_( } metal::float2 _e78 = pos; metal::float2 _e79 = vPos; - float _e83 = params.rule3Distance; + float _e83 = argumentBuffer.params.rule3Distance; if (metal::distance(_e78, _e79) < _e83) { metal::float2 _e85 = cVel; metal::float2 _e86 = vel; @@ -127,18 +130,18 @@ kernel void main_( } metal::float2 _e112 = vVel; metal::float2 _e113 = cMass; - float _e116 = params.rule1Scale; + float _e116 = argumentBuffer.params.rule1Scale; metal::float2 _e119 = colVel; - float _e122 = params.rule2Scale; + float _e122 = argumentBuffer.params.rule2Scale; metal::float2 _e125 = cVel; - float _e128 = params.rule3Scale; + float _e128 = argumentBuffer.params.rule3Scale; vVel = ((_e112 + (_e113 * _e116)) + (_e119 * _e122)) + (_e125 * _e128); metal::float2 _e131 = vVel; metal::float2 _e133 = vVel; vVel = metal::normalize(_e131) * metal::clamp(metal::length(_e133), 0.0, 0.1); metal::float2 _e139 = vPos; metal::float2 _e140 = vVel; - float _e143 = params.deltaT; + float _e143 = argumentBuffer.params.deltaT; vPos = _e139 + (_e140 * _e143); float _e147 = vPos.x; if (_e147 < -1.0) { @@ -157,8 +160,8 @@ kernel void main_( vPos.y = -1.0; } metal::float2 _e174 = vPos; - particlesDst.particles[index].pos = _e174; + argumentBuffer.particlesDst.particles[index].pos = _e174; metal::float2 _e179 = vVel; - particlesDst.particles[index].vel = _e179; + argumentBuffer.particlesDst.particles[index].vel = _e179; return; } diff --git a/tests/out/msl/bounds-check-image-restrict.msl b/tests/out/msl/bounds-check-image-restrict.msl index 9f94ef0a6e..83912918b9 100644 --- a/tests/out/msl/bounds-check-image-restrict.msl +++ b/tests/out/msl/bounds-check-image-restrict.msl @@ -156,27 +156,30 @@ void test_textureStore_3d( struct fragment_shaderOutput { metal::float4 member [[color(0)]]; }; +struct fragment_shaderArgumentBuffer { + metal::texture1d image_1d [[id(0)]]; + metal::texture2d image_2d [[id(1)]]; + metal::texture2d_array image_2d_array [[id(2)]]; + metal::texture3d image_3d [[id(3)]]; + metal::texture2d_ms image_multisampled_2d [[id(4)]]; + metal::texture1d image_storage_1d [[id(5)]]; + metal::texture2d image_storage_2d [[id(6)]]; + metal::texture2d_array image_storage_2d_array [[id(7)]]; + metal::texture3d image_storage_3d [[id(8)]]; +}; fragment fragment_shaderOutput fragment_shader( - metal::texture1d image_1d [[user(fake0)]] -, metal::texture2d image_2d [[user(fake0)]] -, metal::texture2d_array image_2d_array [[user(fake0)]] -, metal::texture3d image_3d [[user(fake0)]] -, metal::texture2d_ms image_multisampled_2d [[user(fake0)]] -, metal::texture1d image_storage_1d [[user(fake0)]] -, metal::texture2d image_storage_2d [[user(fake0)]] -, metal::texture2d_array image_storage_2d_array [[user(fake0)]] -, metal::texture3d image_storage_3d [[user(fake0)]] -) { - metal::float4 _e2 = test_textureLoad_1d(0, 0, image_1d); - metal::float4 _e5 = test_textureLoad_2d(metal::int2 {}, 0, image_2d); - metal::float4 _e9 = test_textureLoad_2d_array_u(metal::int2 {}, 0u, 0, image_2d_array); - metal::float4 _e13 = test_textureLoad_2d_array_s(metal::int2 {}, 0, 0, image_2d_array); - metal::float4 _e16 = test_textureLoad_3d(metal::int3 {}, 0, image_3d); - metal::float4 _e19 = test_textureLoad_multisampled_2d(metal::int2 {}, 0, image_multisampled_2d); - test_textureStore_1d(0, metal::float4 {}, image_storage_1d); - test_textureStore_2d(metal::int2 {}, metal::float4 {}, image_storage_2d); - test_textureStore_2d_array_u(metal::int2 {}, 0u, metal::float4 {}, image_storage_2d_array); - test_textureStore_2d_array_s(metal::int2 {}, 0, metal::float4 {}, image_storage_2d_array); - test_textureStore_3d(metal::int3 {}, metal::float4 {}, image_storage_3d); + device fragment_shaderArgumentBuffer& argumentBuffer [[buffer(0)]] +) { + metal::float4 _e2 = test_textureLoad_1d(0, 0, argumentBuffer.image_1d); + metal::float4 _e5 = test_textureLoad_2d(metal::int2 {}, 0, argumentBuffer.image_2d); + metal::float4 _e9 = test_textureLoad_2d_array_u(metal::int2 {}, 0u, 0, argumentBuffer.image_2d_array); + metal::float4 _e13 = test_textureLoad_2d_array_s(metal::int2 {}, 0, 0, argumentBuffer.image_2d_array); + metal::float4 _e16 = test_textureLoad_3d(metal::int3 {}, 0, argumentBuffer.image_3d); + metal::float4 _e19 = test_textureLoad_multisampled_2d(metal::int2 {}, 0, argumentBuffer.image_multisampled_2d); + test_textureStore_1d(0, metal::float4 {}, argumentBuffer.image_storage_1d); + test_textureStore_2d(metal::int2 {}, metal::float4 {}, argumentBuffer.image_storage_2d); + test_textureStore_2d_array_u(metal::int2 {}, 0u, metal::float4 {}, argumentBuffer.image_storage_2d_array); + test_textureStore_2d_array_s(metal::int2 {}, 0, metal::float4 {}, argumentBuffer.image_storage_2d_array); + test_textureStore_3d(metal::int3 {}, metal::float4 {}, argumentBuffer.image_storage_3d); return fragment_shaderOutput { metal::float4(0.0, 0.0, 0.0, 0.0) }; } diff --git a/tests/out/msl/bounds-check-image-rzsw.msl b/tests/out/msl/bounds-check-image-rzsw.msl index a93014cb27..2a60ef4b7b 100644 --- a/tests/out/msl/bounds-check-image-rzsw.msl +++ b/tests/out/msl/bounds-check-image-rzsw.msl @@ -165,27 +165,30 @@ void test_textureStore_3d( struct fragment_shaderOutput { metal::float4 member [[color(0)]]; }; +struct fragment_shaderArgumentBuffer { + metal::texture1d image_1d [[id(0)]]; + metal::texture2d image_2d [[id(1)]]; + metal::texture2d_array image_2d_array [[id(2)]]; + metal::texture3d image_3d [[id(3)]]; + metal::texture2d_ms image_multisampled_2d [[id(4)]]; + metal::texture1d image_storage_1d [[id(5)]]; + metal::texture2d image_storage_2d [[id(6)]]; + metal::texture2d_array image_storage_2d_array [[id(7)]]; + metal::texture3d image_storage_3d [[id(8)]]; +}; fragment fragment_shaderOutput fragment_shader( - metal::texture1d image_1d [[user(fake0)]] -, metal::texture2d image_2d [[user(fake0)]] -, metal::texture2d_array image_2d_array [[user(fake0)]] -, metal::texture3d image_3d [[user(fake0)]] -, metal::texture2d_ms image_multisampled_2d [[user(fake0)]] -, metal::texture1d image_storage_1d [[user(fake0)]] -, metal::texture2d image_storage_2d [[user(fake0)]] -, metal::texture2d_array image_storage_2d_array [[user(fake0)]] -, metal::texture3d image_storage_3d [[user(fake0)]] -) { - metal::float4 _e2 = test_textureLoad_1d(0, 0, image_1d); - metal::float4 _e5 = test_textureLoad_2d(metal::int2 {}, 0, image_2d); - metal::float4 _e9 = test_textureLoad_2d_array_u(metal::int2 {}, 0u, 0, image_2d_array); - metal::float4 _e13 = test_textureLoad_2d_array_s(metal::int2 {}, 0, 0, image_2d_array); - metal::float4 _e16 = test_textureLoad_3d(metal::int3 {}, 0, image_3d); - metal::float4 _e19 = test_textureLoad_multisampled_2d(metal::int2 {}, 0, image_multisampled_2d); - test_textureStore_1d(0, metal::float4 {}, image_storage_1d); - test_textureStore_2d(metal::int2 {}, metal::float4 {}, image_storage_2d); - test_textureStore_2d_array_u(metal::int2 {}, 0u, metal::float4 {}, image_storage_2d_array); - test_textureStore_2d_array_s(metal::int2 {}, 0, metal::float4 {}, image_storage_2d_array); - test_textureStore_3d(metal::int3 {}, metal::float4 {}, image_storage_3d); + device fragment_shaderArgumentBuffer& argumentBuffer [[buffer(0)]] +) { + metal::float4 _e2 = test_textureLoad_1d(0, 0, argumentBuffer.image_1d); + metal::float4 _e5 = test_textureLoad_2d(metal::int2 {}, 0, argumentBuffer.image_2d); + metal::float4 _e9 = test_textureLoad_2d_array_u(metal::int2 {}, 0u, 0, argumentBuffer.image_2d_array); + metal::float4 _e13 = test_textureLoad_2d_array_s(metal::int2 {}, 0, 0, argumentBuffer.image_2d_array); + metal::float4 _e16 = test_textureLoad_3d(metal::int3 {}, 0, argumentBuffer.image_3d); + metal::float4 _e19 = test_textureLoad_multisampled_2d(metal::int2 {}, 0, argumentBuffer.image_multisampled_2d); + test_textureStore_1d(0, metal::float4 {}, argumentBuffer.image_storage_1d); + test_textureStore_2d(metal::int2 {}, metal::float4 {}, argumentBuffer.image_storage_2d); + test_textureStore_2d_array_u(metal::int2 {}, 0u, metal::float4 {}, argumentBuffer.image_storage_2d_array); + test_textureStore_2d_array_s(metal::int2 {}, 0, metal::float4 {}, argumentBuffer.image_storage_2d_array); + test_textureStore_3d(metal::int3 {}, metal::float4 {}, argumentBuffer.image_storage_3d); return fragment_shaderOutput { metal::float4(0.0, 0.0, 0.0, 0.0) }; } diff --git a/tests/out/msl/collatz.msl b/tests/out/msl/collatz.msl index 88f9521a27..799780c02a 100644 --- a/tests/out/msl/collatz.msl +++ b/tests/out/msl/collatz.msl @@ -45,13 +45,16 @@ uint collatz_iterations( struct main_Input { }; +struct main_ArgumentBuffer { + device PrimeIndices& v_indices [[id(0)]]; +}; kernel void main_( metal::uint3 global_id [[thread_position_in_grid]] -, device PrimeIndices& v_indices [[user(fake0)]] +, device main_ArgumentBuffer& argumentBuffer [[buffer(0)]] , constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] ) { - uint _e9 = v_indices.data[global_id.x]; + uint _e9 = argumentBuffer.v_indices.data[global_id.x]; uint _e10 = collatz_iterations(_e9); - v_indices.data[global_id.x] = _e10; + argumentBuffer.v_indices.data[global_id.x] = _e10; return; } diff --git a/tests/out/msl/globals.msl b/tests/out/msl/globals.msl index d9142c1990..3c43296d2c 100644 --- a/tests/out/msl/globals.msl +++ b/tests/out/msl/globals.msl @@ -59,17 +59,20 @@ void test_msl_packed_vec3_( metal::float3 svm1_ = 2.0 * data.v3_; } +struct main_ArgumentBuffer { + device FooStruct& alignment [[id(0)]]; + device type_6 const& dummy [[id(1)]]; + constant type_8& float_vecs [[id(2)]]; + constant metal::float3& global_vec [[id(22)]]; + constant metal::float3x2& global_mat [[id(23)]]; + constant type_12& global_nested_arrays_of_matrices_2x4_ [[id(24)]]; + constant type_15& global_nested_arrays_of_matrices_4x2_ [[id(26)]]; +}; kernel void main_( metal::uint3 __local_invocation_id [[thread_position_in_threadgroup]] +, device main_ArgumentBuffer& argumentBuffer [[buffer(0)]] , threadgroup type_2& wg , threadgroup metal::atomic_uint& at_1 -, device FooStruct& alignment [[user(fake0)]] -, device type_6 const& dummy [[user(fake0)]] -, constant type_8& float_vecs [[user(fake0)]] -, constant metal::float3& global_vec [[user(fake0)]] -, constant metal::float3x2& global_mat [[user(fake0)]] -, constant type_12& global_nested_arrays_of_matrices_2x4_ [[user(fake0)]] -, constant type_15& global_nested_arrays_of_matrices_4x2_ [[user(fake0)]] , constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] ) { if (metal::all(__local_invocation_id == metal::uint3(0u))) { @@ -79,22 +82,22 @@ kernel void main_( metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); float Foo = {}; bool at = {}; - test_msl_packed_vec3_(alignment); - metal::float4x2 _e8 = global_nested_arrays_of_matrices_4x2_.inner[0].inner[0]; - metal::float4 _e16 = global_nested_arrays_of_matrices_2x4_.inner[0].inner[0][0]; + test_msl_packed_vec3_(argumentBuffer.alignment); + metal::float4x2 _e8 = argumentBuffer.global_nested_arrays_of_matrices_4x2_.inner[0].inner[0]; + metal::float4 _e16 = argumentBuffer.global_nested_arrays_of_matrices_2x4_.inner[0].inner[0][0]; wg.inner[7] = (_e8 * _e16).x; - metal::float3x2 _e23 = global_mat; - metal::float3 _e25 = global_vec; + metal::float3x2 _e23 = argumentBuffer.global_mat; + metal::float3 _e25 = argumentBuffer.global_vec; wg.inner[6] = (_e23 * _e25).x; - float _e35 = dummy[1].y; + float _e35 = argumentBuffer.dummy[1].y; wg.inner[5] = _e35; - float _e43 = float_vecs.inner[0].w; + float _e43 = argumentBuffer.float_vecs.inner[0].w; wg.inner[4] = _e43; - float _e49 = alignment.v1_; + float _e49 = argumentBuffer.alignment.v1_; wg.inner[3] = _e49; - float _e56 = alignment.v3_[0]; + float _e56 = argumentBuffer.alignment.v3_[0]; wg.inner[2] = _e56; - alignment.v1_ = 4.0; + argumentBuffer.alignment.v1_ = 4.0; wg.inner[1] = static_cast(1 + (_buffer_sizes.size3 - 0 - 8) / 8); metal::atomic_store_explicit(&at_1, 2u, metal::memory_order_relaxed); Foo = 1.0; diff --git a/tests/out/msl/image.msl b/tests/out/msl/image.msl index e390c2e0fc..fece041428 100644 --- a/tests/out/msl/image.msl +++ b/tests/out/msl/image.msl @@ -7,47 +7,53 @@ using metal::uint; struct main_Input { }; +struct main_ArgumentBuffer { + metal::texture2d image_mipmapped_src [[id(0)]]; + metal::texture2d_ms image_multisampled_src [[id(1)]]; + metal::texture2d image_storage_src [[id(2)]]; + metal::texture2d_array image_array_src [[id(3)]]; + metal::texture1d image_1d_src [[id(4)]]; + metal::texture1d image_dst [[id(5)]]; +}; kernel void main_( metal::uint3 local_id [[thread_position_in_threadgroup]] -, metal::texture2d image_mipmapped_src [[user(fake0)]] -, metal::texture2d_ms image_multisampled_src [[user(fake0)]] -, metal::texture2d image_storage_src [[user(fake0)]] -, metal::texture2d_array image_array_src [[user(fake0)]] -, metal::texture1d image_1d_src [[user(fake0)]] -, metal::texture1d image_dst [[user(fake0)]] +, device main_ArgumentBuffer& argumentBuffer [[buffer(0)]] ) { - metal::uint2 dim = metal::uint2(image_storage_src.get_width(), image_storage_src.get_height()); + metal::uint2 dim = metal::uint2(argumentBuffer.image_storage_src.get_width(), argumentBuffer.image_storage_src.get_height()); metal::int2 itc = static_cast(dim * local_id.xy) % metal::int2(10, 20); - metal::uint4 value1_ = image_mipmapped_src.read(metal::uint2(itc), static_cast(local_id.z)); - metal::uint4 value2_ = image_multisampled_src.read(metal::uint2(itc), static_cast(local_id.z)); - metal::uint4 value4_ = image_storage_src.read(metal::uint2(itc)); - metal::uint4 value5_ = image_array_src.read(metal::uint2(itc), local_id.z, static_cast(local_id.z) + 1); - metal::uint4 value6_ = image_array_src.read(metal::uint2(itc), static_cast(local_id.z), static_cast(local_id.z) + 1); - metal::uint4 value7_ = image_1d_src.read(uint(static_cast(local_id.x))); - metal::uint4 value1u = image_mipmapped_src.read(metal::uint2(static_cast(itc)), static_cast(local_id.z)); - metal::uint4 value2u = image_multisampled_src.read(metal::uint2(static_cast(itc)), static_cast(local_id.z)); - metal::uint4 value4u = image_storage_src.read(metal::uint2(static_cast(itc))); - metal::uint4 value5u = image_array_src.read(metal::uint2(static_cast(itc)), local_id.z, static_cast(local_id.z) + 1); - metal::uint4 value6u = image_array_src.read(metal::uint2(static_cast(itc)), static_cast(local_id.z), static_cast(local_id.z) + 1); - metal::uint4 value7u = image_1d_src.read(uint(static_cast(local_id.x))); - image_dst.write((((value1_ + value2_) + value4_) + value5_) + value6_, uint(itc.x)); - image_dst.write((((value1u + value2u) + value4u) + value5u) + value6u, uint(static_cast(itc.x))); + metal::uint4 value1_ = argumentBuffer.image_mipmapped_src.read(metal::uint2(itc), static_cast(local_id.z)); + metal::uint4 value2_ = argumentBuffer.image_multisampled_src.read(metal::uint2(itc), static_cast(local_id.z)); + metal::uint4 value4_ = argumentBuffer.image_storage_src.read(metal::uint2(itc)); + metal::uint4 value5_ = argumentBuffer.image_array_src.read(metal::uint2(itc), local_id.z, static_cast(local_id.z) + 1); + metal::uint4 value6_ = argumentBuffer.image_array_src.read(metal::uint2(itc), static_cast(local_id.z), static_cast(local_id.z) + 1); + metal::uint4 value7_ = argumentBuffer.image_1d_src.read(uint(static_cast(local_id.x))); + metal::uint4 value1u = argumentBuffer.image_mipmapped_src.read(metal::uint2(static_cast(itc)), static_cast(local_id.z)); + metal::uint4 value2u = argumentBuffer.image_multisampled_src.read(metal::uint2(static_cast(itc)), static_cast(local_id.z)); + metal::uint4 value4u = argumentBuffer.image_storage_src.read(metal::uint2(static_cast(itc))); + metal::uint4 value5u = argumentBuffer.image_array_src.read(metal::uint2(static_cast(itc)), local_id.z, static_cast(local_id.z) + 1); + metal::uint4 value6u = argumentBuffer.image_array_src.read(metal::uint2(static_cast(itc)), static_cast(local_id.z), static_cast(local_id.z) + 1); + metal::uint4 value7u = argumentBuffer.image_1d_src.read(uint(static_cast(local_id.x))); + argumentBuffer.image_dst.write((((value1_ + value2_) + value4_) + value5_) + value6_, uint(itc.x)); + argumentBuffer.image_dst.write((((value1u + value2u) + value4u) + value5u) + value6u, uint(static_cast(itc.x))); return; } struct depth_loadInput { }; +struct depth_loadArgumentBuffer { + metal::depth2d_ms image_depth_multisampled_src [[id(0)]]; + metal::texture2d image_storage_src [[id(1)]]; + metal::texture1d image_dst [[id(2)]]; +}; kernel void depth_load( metal::uint3 local_id_1 [[thread_position_in_threadgroup]] -, metal::depth2d_ms image_depth_multisampled_src [[user(fake0)]] -, metal::texture2d image_storage_src [[user(fake0)]] -, metal::texture1d image_dst [[user(fake0)]] +, device depth_loadArgumentBuffer& argumentBuffer [[buffer(0)]] ) { - metal::uint2 dim_1 = metal::uint2(image_storage_src.get_width(), image_storage_src.get_height()); + metal::uint2 dim_1 = metal::uint2(argumentBuffer.image_storage_src.get_width(), argumentBuffer.image_storage_src.get_height()); metal::int2 itc_1 = static_cast(dim_1 * local_id_1.xy) % metal::int2(10, 20); - float val = image_depth_multisampled_src.read(metal::uint2(itc_1), static_cast(local_id_1.z)); - image_dst.write(metal::uint4(static_cast(val)), uint(itc_1.x)); + float val = argumentBuffer.image_depth_multisampled_src.read(metal::uint2(itc_1), static_cast(local_id_1.z)); + argumentBuffer.image_dst.write(metal::uint4(static_cast(val)), uint(itc_1.x)); return; } @@ -55,28 +61,31 @@ kernel void depth_load( struct queriesOutput { metal::float4 member_2 [[position]]; }; +struct queriesArgumentBuffer { + metal::texture1d image_1d [[id(0)]]; + metal::texture2d image_2d [[id(1)]]; + metal::texture2d_array image_2d_array [[id(2)]]; + metal::texturecube image_cube [[id(3)]]; + metal::texturecube_array image_cube_array [[id(4)]]; + metal::texture3d image_3d [[id(5)]]; + metal::texture2d_ms image_aa [[id(6)]]; +}; vertex queriesOutput queries( - metal::texture1d image_1d [[user(fake0)]] -, metal::texture2d image_2d [[user(fake0)]] -, metal::texture2d_array image_2d_array [[user(fake0)]] -, metal::texturecube image_cube [[user(fake0)]] -, metal::texturecube_array image_cube_array [[user(fake0)]] -, metal::texture3d image_3d [[user(fake0)]] -, metal::texture2d_ms image_aa [[user(fake0)]] + device queriesArgumentBuffer& argumentBuffer [[buffer(0)]] ) { - uint dim_1d = image_1d.get_width(); - uint dim_1d_lod = image_1d.get_width(); - metal::uint2 dim_2d = metal::uint2(image_2d.get_width(), image_2d.get_height()); - metal::uint2 dim_2d_lod = metal::uint2(image_2d.get_width(1), image_2d.get_height(1)); - metal::uint2 dim_2d_array = metal::uint2(image_2d_array.get_width(), image_2d_array.get_height()); - metal::uint2 dim_2d_array_lod = metal::uint2(image_2d_array.get_width(1), image_2d_array.get_height(1)); - metal::uint2 dim_cube = metal::uint2(image_cube.get_width()); - metal::uint2 dim_cube_lod = metal::uint2(image_cube.get_width(1)); - metal::uint2 dim_cube_array = metal::uint2(image_cube_array.get_width()); - metal::uint2 dim_cube_array_lod = metal::uint2(image_cube_array.get_width(1)); - metal::uint3 dim_3d = metal::uint3(image_3d.get_width(), image_3d.get_height(), image_3d.get_depth()); - metal::uint3 dim_3d_lod = metal::uint3(image_3d.get_width(1), image_3d.get_height(1), image_3d.get_depth(1)); - metal::uint2 dim_2s_ms = metal::uint2(image_aa.get_width(), image_aa.get_height()); + uint dim_1d = argumentBuffer.image_1d.get_width(); + uint dim_1d_lod = argumentBuffer.image_1d.get_width(); + metal::uint2 dim_2d = metal::uint2(argumentBuffer.image_2d.get_width(), argumentBuffer.image_2d.get_height()); + metal::uint2 dim_2d_lod = metal::uint2(argumentBuffer.image_2d.get_width(1), argumentBuffer.image_2d.get_height(1)); + metal::uint2 dim_2d_array = metal::uint2(argumentBuffer.image_2d_array.get_width(), argumentBuffer.image_2d_array.get_height()); + metal::uint2 dim_2d_array_lod = metal::uint2(argumentBuffer.image_2d_array.get_width(1), argumentBuffer.image_2d_array.get_height(1)); + metal::uint2 dim_cube = metal::uint2(argumentBuffer.image_cube.get_width()); + metal::uint2 dim_cube_lod = metal::uint2(argumentBuffer.image_cube.get_width(1)); + metal::uint2 dim_cube_array = metal::uint2(argumentBuffer.image_cube_array.get_width()); + metal::uint2 dim_cube_array_lod = metal::uint2(argumentBuffer.image_cube_array.get_width(1)); + metal::uint3 dim_3d = metal::uint3(argumentBuffer.image_3d.get_width(), argumentBuffer.image_3d.get_height(), argumentBuffer.image_3d.get_depth()); + metal::uint3 dim_3d_lod = metal::uint3(argumentBuffer.image_3d.get_width(1), argumentBuffer.image_3d.get_height(1), argumentBuffer.image_3d.get_depth(1)); + metal::uint2 dim_2s_ms = metal::uint2(argumentBuffer.image_aa.get_width(), argumentBuffer.image_aa.get_height()); uint sum = (((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z; return queriesOutput { metal::float4(static_cast(sum)) }; } @@ -85,22 +94,25 @@ vertex queriesOutput queries( struct levels_queriesOutput { metal::float4 member_3 [[position]]; }; +struct levels_queriesArgumentBuffer { + metal::texture2d image_2d [[id(0)]]; + metal::texture2d_array image_2d_array [[id(1)]]; + metal::texturecube image_cube [[id(2)]]; + metal::texturecube_array image_cube_array [[id(3)]]; + metal::texture3d image_3d [[id(4)]]; + metal::texture2d_ms image_aa [[id(5)]]; +}; vertex levels_queriesOutput levels_queries( - metal::texture2d image_2d [[user(fake0)]] -, metal::texture2d_array image_2d_array [[user(fake0)]] -, metal::texturecube image_cube [[user(fake0)]] -, metal::texturecube_array image_cube_array [[user(fake0)]] -, metal::texture3d image_3d [[user(fake0)]] -, metal::texture2d_ms image_aa [[user(fake0)]] + device levels_queriesArgumentBuffer& argumentBuffer [[buffer(0)]] ) { - uint num_levels_2d = image_2d.get_num_mip_levels(); - uint num_levels_2d_array = image_2d_array.get_num_mip_levels(); - uint num_layers_2d = image_2d_array.get_array_size(); - uint num_levels_cube = image_cube.get_num_mip_levels(); - uint num_levels_cube_array = image_cube_array.get_num_mip_levels(); - uint num_layers_cube = image_cube_array.get_array_size(); - uint num_levels_3d = image_3d.get_num_mip_levels(); - uint num_samples_aa = image_aa.get_num_samples(); + uint num_levels_2d = argumentBuffer.image_2d.get_num_mip_levels(); + uint num_levels_2d_array = argumentBuffer.image_2d_array.get_num_mip_levels(); + uint num_layers_2d = argumentBuffer.image_2d_array.get_array_size(); + uint num_levels_cube = argumentBuffer.image_cube.get_num_mip_levels(); + uint num_levels_cube_array = argumentBuffer.image_cube_array.get_num_mip_levels(); + uint num_layers_cube = argumentBuffer.image_cube_array.get_array_size(); + uint num_levels_3d = argumentBuffer.image_3d.get_num_mip_levels(); + uint num_samples_aa = argumentBuffer.image_aa.get_num_samples(); uint sum_1 = ((((((num_layers_2d + num_layers_cube) + num_samples_aa) + num_levels_2d) + num_levels_2d_array) + num_levels_3d) + num_levels_cube) + num_levels_cube_array; return levels_queriesOutput { metal::float4(static_cast(sum_1)) }; } @@ -109,80 +121,83 @@ vertex levels_queriesOutput levels_queries( struct texture_sampleOutput { metal::float4 member_4 [[color(0)]]; }; +struct texture_sampleArgumentBuffer { + metal::texture1d image_1d [[id(0)]]; + metal::texture2d image_2d [[id(1)]]; + metal::texture2d_array image_2d_array [[id(2)]]; + metal::texturecube_array image_cube_array [[id(3)]]; + metal::sampler sampler_reg [[id(4)]]; +}; fragment texture_sampleOutput texture_sample( - metal::texture1d image_1d [[user(fake0)]] -, metal::texture2d image_2d [[user(fake0)]] -, metal::texture2d_array image_2d_array [[user(fake0)]] -, metal::texturecube_array image_cube_array [[user(fake0)]] -, metal::sampler sampler_reg [[user(fake0)]] + device texture_sampleArgumentBuffer& argumentBuffer [[buffer(0)]] ) { metal::float4 a = {}; metal::float2 tc = metal::float2(0.5); metal::float3 tc3_ = metal::float3(0.5); - metal::float4 _e9 = image_1d.sample(sampler_reg, tc.x); + metal::float4 _e9 = argumentBuffer.image_1d.sample(argumentBuffer.sampler_reg, tc.x); metal::float4 _e10 = a; a = _e10 + _e9; - metal::float4 _e14 = image_2d.sample(sampler_reg, tc); + metal::float4 _e14 = argumentBuffer.image_2d.sample(argumentBuffer.sampler_reg, tc); metal::float4 _e15 = a; a = _e15 + _e14; - metal::float4 _e19 = image_2d.sample(sampler_reg, tc, metal::int2(3, 1)); + metal::float4 _e19 = argumentBuffer.image_2d.sample(argumentBuffer.sampler_reg, tc, metal::int2(3, 1)); metal::float4 _e20 = a; a = _e20 + _e19; - metal::float4 _e24 = image_2d.sample(sampler_reg, tc, metal::level(2.3)); + metal::float4 _e24 = argumentBuffer.image_2d.sample(argumentBuffer.sampler_reg, tc, metal::level(2.3)); metal::float4 _e25 = a; a = _e25 + _e24; - metal::float4 _e29 = image_2d.sample(sampler_reg, tc, metal::level(2.3), metal::int2(3, 1)); + metal::float4 _e29 = argumentBuffer.image_2d.sample(argumentBuffer.sampler_reg, tc, metal::level(2.3), metal::int2(3, 1)); metal::float4 _e30 = a; a = _e30 + _e29; - metal::float4 _e35 = image_2d.sample(sampler_reg, tc, metal::bias(2.0), metal::int2(3, 1)); + metal::float4 _e35 = argumentBuffer.image_2d.sample(argumentBuffer.sampler_reg, tc, metal::bias(2.0), metal::int2(3, 1)); metal::float4 _e36 = a; a = _e36 + _e35; - metal::float4 _e41 = image_2d_array.sample(sampler_reg, tc, 0u); + metal::float4 _e41 = argumentBuffer.image_2d_array.sample(argumentBuffer.sampler_reg, tc, 0u); metal::float4 _e42 = a; a = _e42 + _e41; - metal::float4 _e47 = image_2d_array.sample(sampler_reg, tc, 0u, metal::int2(3, 1)); + metal::float4 _e47 = argumentBuffer.image_2d_array.sample(argumentBuffer.sampler_reg, tc, 0u, metal::int2(3, 1)); metal::float4 _e48 = a; a = _e48 + _e47; - metal::float4 _e53 = image_2d_array.sample(sampler_reg, tc, 0u, metal::level(2.3)); + metal::float4 _e53 = argumentBuffer.image_2d_array.sample(argumentBuffer.sampler_reg, tc, 0u, metal::level(2.3)); metal::float4 _e54 = a; a = _e54 + _e53; - metal::float4 _e59 = image_2d_array.sample(sampler_reg, tc, 0u, metal::level(2.3), metal::int2(3, 1)); + metal::float4 _e59 = argumentBuffer.image_2d_array.sample(argumentBuffer.sampler_reg, tc, 0u, metal::level(2.3), metal::int2(3, 1)); metal::float4 _e60 = a; a = _e60 + _e59; - metal::float4 _e66 = image_2d_array.sample(sampler_reg, tc, 0u, metal::bias(2.0), metal::int2(3, 1)); + metal::float4 _e66 = argumentBuffer.image_2d_array.sample(argumentBuffer.sampler_reg, tc, 0u, metal::bias(2.0), metal::int2(3, 1)); metal::float4 _e67 = a; a = _e67 + _e66; - metal::float4 _e72 = image_2d_array.sample(sampler_reg, tc, 0); + metal::float4 _e72 = argumentBuffer.image_2d_array.sample(argumentBuffer.sampler_reg, tc, 0); metal::float4 _e73 = a; a = _e73 + _e72; - metal::float4 _e78 = image_2d_array.sample(sampler_reg, tc, 0, metal::int2(3, 1)); + metal::float4 _e78 = argumentBuffer.image_2d_array.sample(argumentBuffer.sampler_reg, tc, 0, metal::int2(3, 1)); metal::float4 _e79 = a; a = _e79 + _e78; - metal::float4 _e84 = image_2d_array.sample(sampler_reg, tc, 0, metal::level(2.3)); + metal::float4 _e84 = argumentBuffer.image_2d_array.sample(argumentBuffer.sampler_reg, tc, 0, metal::level(2.3)); metal::float4 _e85 = a; a = _e85 + _e84; - metal::float4 _e90 = image_2d_array.sample(sampler_reg, tc, 0, metal::level(2.3), metal::int2(3, 1)); + metal::float4 _e90 = argumentBuffer.image_2d_array.sample(argumentBuffer.sampler_reg, tc, 0, metal::level(2.3), metal::int2(3, 1)); metal::float4 _e91 = a; a = _e91 + _e90; - metal::float4 _e97 = image_2d_array.sample(sampler_reg, tc, 0, metal::bias(2.0), metal::int2(3, 1)); + metal::float4 _e97 = argumentBuffer.image_2d_array.sample(argumentBuffer.sampler_reg, tc, 0, metal::bias(2.0), metal::int2(3, 1)); metal::float4 _e98 = a; a = _e98 + _e97; - metal::float4 _e103 = image_cube_array.sample(sampler_reg, tc3_, 0u); + metal::float4 _e103 = argumentBuffer.image_cube_array.sample(argumentBuffer.sampler_reg, tc3_, 0u); metal::float4 _e104 = a; a = _e104 + _e103; - metal::float4 _e109 = image_cube_array.sample(sampler_reg, tc3_, 0u, metal::level(2.3)); + metal::float4 _e109 = argumentBuffer.image_cube_array.sample(argumentBuffer.sampler_reg, tc3_, 0u, metal::level(2.3)); metal::float4 _e110 = a; a = _e110 + _e109; - metal::float4 _e116 = image_cube_array.sample(sampler_reg, tc3_, 0u, metal::bias(2.0)); + metal::float4 _e116 = argumentBuffer.image_cube_array.sample(argumentBuffer.sampler_reg, tc3_, 0u, metal::bias(2.0)); metal::float4 _e117 = a; a = _e117 + _e116; - metal::float4 _e122 = image_cube_array.sample(sampler_reg, tc3_, 0); + metal::float4 _e122 = argumentBuffer.image_cube_array.sample(argumentBuffer.sampler_reg, tc3_, 0); metal::float4 _e123 = a; a = _e123 + _e122; - metal::float4 _e128 = image_cube_array.sample(sampler_reg, tc3_, 0, metal::level(2.3)); + metal::float4 _e128 = argumentBuffer.image_cube_array.sample(argumentBuffer.sampler_reg, tc3_, 0, metal::level(2.3)); metal::float4 _e129 = a; a = _e129 + _e128; - metal::float4 _e135 = image_cube_array.sample(sampler_reg, tc3_, 0, metal::bias(2.0)); + metal::float4 _e135 = argumentBuffer.image_cube_array.sample(argumentBuffer.sampler_reg, tc3_, 0, metal::bias(2.0)); metal::float4 _e136 = a; a = _e136 + _e135; metal::float4 _e138 = a; @@ -193,37 +208,40 @@ fragment texture_sampleOutput texture_sample( struct texture_sample_comparisonOutput { float member_5 [[color(0)]]; }; +struct texture_sample_comparisonArgumentBuffer { + metal::sampler sampler_cmp [[id(0)]]; + metal::depth2d image_2d_depth [[id(1)]]; + metal::depth2d_array image_2d_array_depth [[id(2)]]; + metal::depthcube image_cube_depth [[id(3)]]; +}; fragment texture_sample_comparisonOutput texture_sample_comparison( - metal::sampler sampler_cmp [[user(fake0)]] -, metal::depth2d image_2d_depth [[user(fake0)]] -, metal::depth2d_array image_2d_array_depth [[user(fake0)]] -, metal::depthcube image_cube_depth [[user(fake0)]] + device texture_sample_comparisonArgumentBuffer& argumentBuffer [[buffer(0)]] ) { float a_1 = {}; metal::float2 tc_1 = metal::float2(0.5); metal::float3 tc3_1 = metal::float3(0.5); - float _e8 = image_2d_depth.sample_compare(sampler_cmp, tc_1, 0.5); + float _e8 = argumentBuffer.image_2d_depth.sample_compare(argumentBuffer.sampler_cmp, tc_1, 0.5); float _e9 = a_1; a_1 = _e9 + _e8; - float _e14 = image_2d_array_depth.sample_compare(sampler_cmp, tc_1, 0u, 0.5); + float _e14 = argumentBuffer.image_2d_array_depth.sample_compare(argumentBuffer.sampler_cmp, tc_1, 0u, 0.5); float _e15 = a_1; a_1 = _e15 + _e14; - float _e20 = image_2d_array_depth.sample_compare(sampler_cmp, tc_1, 0, 0.5); + float _e20 = argumentBuffer.image_2d_array_depth.sample_compare(argumentBuffer.sampler_cmp, tc_1, 0, 0.5); float _e21 = a_1; a_1 = _e21 + _e20; - float _e25 = image_cube_depth.sample_compare(sampler_cmp, tc3_1, 0.5); + float _e25 = argumentBuffer.image_cube_depth.sample_compare(argumentBuffer.sampler_cmp, tc3_1, 0.5); float _e26 = a_1; a_1 = _e26 + _e25; - float _e30 = image_2d_depth.sample_compare(sampler_cmp, tc_1, 0.5); + float _e30 = argumentBuffer.image_2d_depth.sample_compare(argumentBuffer.sampler_cmp, tc_1, 0.5); float _e31 = a_1; a_1 = _e31 + _e30; - float _e36 = image_2d_array_depth.sample_compare(sampler_cmp, tc_1, 0u, 0.5); + float _e36 = argumentBuffer.image_2d_array_depth.sample_compare(argumentBuffer.sampler_cmp, tc_1, 0u, 0.5); float _e37 = a_1; a_1 = _e37 + _e36; - float _e42 = image_2d_array_depth.sample_compare(sampler_cmp, tc_1, 0, 0.5); + float _e42 = argumentBuffer.image_2d_array_depth.sample_compare(argumentBuffer.sampler_cmp, tc_1, 0, 0.5); float _e43 = a_1; a_1 = _e43 + _e42; - float _e47 = image_cube_depth.sample_compare(sampler_cmp, tc3_1, 0.5); + float _e47 = argumentBuffer.image_cube_depth.sample_compare(argumentBuffer.sampler_cmp, tc3_1, 0.5); float _e48 = a_1; a_1 = _e48 + _e47; float _e50 = a_1; @@ -234,21 +252,24 @@ fragment texture_sample_comparisonOutput texture_sample_comparison( struct gatherOutput { metal::float4 member_6 [[color(0)]]; }; +struct gatherArgumentBuffer { + metal::texture2d image_2d [[id(0)]]; + metal::texture2d image_2d_u32_ [[id(1)]]; + metal::texture2d image_2d_i32_ [[id(2)]]; + metal::sampler sampler_reg [[id(3)]]; + metal::sampler sampler_cmp [[id(4)]]; + metal::depth2d image_2d_depth [[id(5)]]; +}; fragment gatherOutput gather( - metal::texture2d image_2d [[user(fake0)]] -, metal::texture2d image_2d_u32_ [[user(fake0)]] -, metal::texture2d image_2d_i32_ [[user(fake0)]] -, metal::sampler sampler_reg [[user(fake0)]] -, metal::sampler sampler_cmp [[user(fake0)]] -, metal::depth2d image_2d_depth [[user(fake0)]] + device gatherArgumentBuffer& argumentBuffer [[buffer(0)]] ) { metal::float2 tc_2 = metal::float2(0.5); - metal::float4 s2d = image_2d.gather(sampler_reg, tc_2, metal::int2(0), metal::component::y); - metal::float4 s2d_offset = image_2d.gather(sampler_reg, tc_2, metal::int2(3, 1), metal::component::w); - metal::float4 s2d_depth = image_2d_depth.gather_compare(sampler_cmp, tc_2, 0.5); - metal::float4 s2d_depth_offset = image_2d_depth.gather_compare(sampler_cmp, tc_2, 0.5, metal::int2(3, 1)); - metal::uint4 u = image_2d_u32_.gather(sampler_reg, tc_2); - metal::int4 i = image_2d_i32_.gather(sampler_reg, tc_2); + metal::float4 s2d = argumentBuffer.image_2d.gather(argumentBuffer.sampler_reg, tc_2, metal::int2(0), metal::component::y); + metal::float4 s2d_offset = argumentBuffer.image_2d.gather(argumentBuffer.sampler_reg, tc_2, metal::int2(3, 1), metal::component::w); + metal::float4 s2d_depth = argumentBuffer.image_2d_depth.gather_compare(argumentBuffer.sampler_cmp, tc_2, 0.5); + metal::float4 s2d_depth_offset = argumentBuffer.image_2d_depth.gather_compare(argumentBuffer.sampler_cmp, tc_2, 0.5, metal::int2(3, 1)); + metal::uint4 u = argumentBuffer.image_2d_u32_.gather(argumentBuffer.sampler_reg, tc_2); + metal::int4 i = argumentBuffer.image_2d_i32_.gather(argumentBuffer.sampler_reg, tc_2); metal::float4 f = static_cast(u) + static_cast(i); return gatherOutput { (((s2d + s2d_offset) + s2d_depth) + s2d_depth_offset) + f }; } @@ -257,12 +278,15 @@ fragment gatherOutput gather( struct depth_no_comparisonOutput { metal::float4 member_7 [[color(0)]]; }; +struct depth_no_comparisonArgumentBuffer { + metal::sampler sampler_reg [[id(0)]]; + metal::depth2d image_2d_depth [[id(1)]]; +}; fragment depth_no_comparisonOutput depth_no_comparison( - metal::sampler sampler_reg [[user(fake0)]] -, metal::depth2d image_2d_depth [[user(fake0)]] + device depth_no_comparisonArgumentBuffer& argumentBuffer [[buffer(0)]] ) { metal::float2 tc_3 = metal::float2(0.5); - float s2d_1 = image_2d_depth.sample(sampler_reg, tc_3); - metal::float4 s2d_gather = image_2d_depth.gather(sampler_reg, tc_3); + float s2d_1 = argumentBuffer.image_2d_depth.sample(argumentBuffer.sampler_reg, tc_3); + metal::float4 s2d_gather = argumentBuffer.image_2d_depth.gather(argumentBuffer.sampler_reg, tc_3); return depth_no_comparisonOutput { metal::float4(s2d_1) + s2d_gather }; } diff --git a/tests/out/msl/padding.msl b/tests/out/msl/padding.msl index 4d99bb4c4c..b9323de240 100644 --- a/tests/out/msl/padding.msl +++ b/tests/out/msl/padding.msl @@ -26,13 +26,16 @@ struct Test3_ { struct vertex_Output { metal::float4 member [[position]]; }; +struct vertex_ArgumentBuffer { + constant Test& input1_ [[id(0)]]; + constant Test2_& input2_ [[id(1)]]; + constant Test3_& input3_ [[id(2)]]; +}; vertex vertex_Output vertex_( - constant Test& input1_ [[buffer(0)]] -, constant Test2_& input2_ [[buffer(1)]] -, constant Test3_& input3_ [[buffer(2)]] + device vertex_ArgumentBuffer& argumentBuffer [[buffer(0)]] ) { - float _e4 = input1_.b; - float _e8 = input2_.b; - float _e12 = input3_.b; + float _e4 = argumentBuffer.input1_.b; + float _e8 = argumentBuffer.input2_.b; + float _e12 = argumentBuffer.input3_.b; return vertex_Output { ((metal::float4(1.0) * _e4) * _e8) * _e12 }; } diff --git a/tests/out/msl/quad.msl b/tests/out/msl/quad.msl index 5fa5788aac..a08c618e6f 100644 --- a/tests/out/msl/quad.msl +++ b/tests/out/msl/quad.msl @@ -34,13 +34,16 @@ struct frag_mainInput { struct frag_mainOutput { metal::float4 member_1 [[color(0)]]; }; +struct frag_mainArgumentBuffer { + metal::texture2d u_texture [[id(0)]]; + metal::sampler u_sampler [[id(1)]]; +}; fragment frag_mainOutput frag_main( frag_mainInput varyings_1 [[stage_in]] -, metal::texture2d u_texture [[user(fake0)]] -, metal::sampler u_sampler [[user(fake0)]] +, device frag_mainArgumentBuffer& argumentBuffer [[buffer(0)]] ) { const auto uv_1 = varyings_1.uv_1; - metal::float4 color = u_texture.sample(u_sampler, uv_1); + metal::float4 color = argumentBuffer.u_texture.sample(argumentBuffer.u_sampler, uv_1); if (color.w == 0.0) { metal::discard_fragment(); } diff --git a/tests/out/msl/ray-query.msl b/tests/out/msl/ray-query.msl index 0d4560f313..dfbcba2e65 100644 --- a/tests/out/msl/ray-query.msl +++ b/tests/out/msl/ray-query.msl @@ -51,9 +51,12 @@ metal::float3 get_torus_normal( return metal::normalize(world_point - world_point_on_guiding_line); } +struct main_ArgumentBuffer { + metal::raytracing::instance_acceleration_structure acc_struct [[id(0)]]; + device Output& output [[id(1)]]; +}; kernel void main_( - metal::raytracing::instance_acceleration_structure acc_struct [[user(fake0)]] -, device Output& output [[user(fake0)]] + device main_ArgumentBuffer& argumentBuffer [[buffer(0)]] ) { _RayQuery rq = {}; metal::float3 dir = metal::float3(0.0, 1.0, 0.0); @@ -62,7 +65,7 @@ kernel void main_( rq.intersector.set_opacity_cull_mode((_e12.flags & 64) != 0 ? metal::raytracing::opacity_cull_mode::opaque : (_e12.flags & 128) != 0 ? metal::raytracing::opacity_cull_mode::non_opaque : metal::raytracing::opacity_cull_mode::none); rq.intersector.force_opacity((_e12.flags & 1) != 0 ? metal::raytracing::forced_opacity::opaque : (_e12.flags & 2) != 0 ? metal::raytracing::forced_opacity::non_opaque : metal::raytracing::forced_opacity::none); rq.intersector.accept_any_intersection((_e12.flags & 4) != 0); - rq.intersection = rq.intersector.intersect(metal::raytracing::ray(_e12.origin, _e12.dir, _e12.tmin, _e12.tmax), acc_struct, _e12.cull_mask); rq.ready = true; + rq.intersection = rq.intersector.intersect(metal::raytracing::ray(_e12.origin, _e12.dir, _e12.tmin, _e12.tmax), argumentBuffer.acc_struct, _e12.cull_mask); rq.ready = true; while(true) { bool _e13 = rq.ready; rq.ready = false; @@ -72,8 +75,8 @@ kernel void main_( } } RayIntersection intersection_1 = RayIntersection {_map_intersection_type(rq.intersection.type), rq.intersection.distance, rq.intersection.user_instance_id, rq.intersection.instance_id, {}, rq.intersection.geometry_id, rq.intersection.primitive_id, rq.intersection.triangle_barycentric_coord, rq.intersection.triangle_front_facing, {}, rq.intersection.object_to_world_transform, rq.intersection.world_to_object_transform}; - output.visible_ = static_cast(intersection_1.kind == 0u); + argumentBuffer.output.visible_ = static_cast(intersection_1.kind == 0u); metal::float3 _e25 = get_torus_normal(dir * intersection_1.t, intersection_1); - output.normal = _e25; + argumentBuffer.output.normal = _e25; return; } diff --git a/tests/out/msl/resource-binding-map.msl b/tests/out/msl/resource-binding-map.msl index b4a53d97b5..a40d2395eb 100644 --- a/tests/out/msl/resource-binding-map.msl +++ b/tests/out/msl/resource-binding-map.msl @@ -16,9 +16,13 @@ struct entry_point_oneInput { struct entry_point_oneOutput { metal::float4 member [[color(0)]]; }; +struct entry_point_oneArgumentBuffer { + metal::texture2d t1_ [[id(0)]]; + metal::sampler s1_ [[id(1)]]; +}; fragment entry_point_oneOutput entry_point_one( metal::float4 pos [[position]] -, metal::texture2d t1_ [[texture(0)]] +, device entry_point_oneArgumentBuffer& argumentBuffer [[buffer(0)]] ) { constexpr metal::sampler s1_( metal::s_address::clamp_to_edge, @@ -28,7 +32,7 @@ fragment entry_point_oneOutput entry_point_one( metal::min_filter::linear, metal::coord::normalized ); - metal::float4 _e4 = t1_.sample(s1_, pos.xy); + metal::float4 _e4 = argumentBuffer.t1_.sample(argumentBuffer.s1_, pos.xy); return entry_point_oneOutput { _e4 }; } @@ -36,13 +40,16 @@ fragment entry_point_oneOutput entry_point_one( struct entry_point_twoOutput { metal::float4 member_1 [[color(0)]]; }; +struct entry_point_twoArgumentBuffer { + metal::texture2d t1_ [[id(0)]]; + metal::sampler s1_ [[id(1)]]; + constant metal::float2& uniformOne [[id(2)]]; +}; fragment entry_point_twoOutput entry_point_two( - metal::texture2d t1_ [[texture(0)]] -, metal::sampler s1_ [[sampler(0)]] -, constant metal::float2& uniformOne [[buffer(0)]] + device entry_point_twoArgumentBuffer& argumentBuffer [[buffer(0)]] ) { - metal::float2 _e3 = uniformOne; - metal::float4 _e4 = t1_.sample(s1_, _e3); + metal::float2 _e3 = argumentBuffer.uniformOne; + metal::float4 _e4 = argumentBuffer.t1_.sample(argumentBuffer.s1_, _e3); return entry_point_twoOutput { _e4 }; } @@ -50,12 +57,16 @@ fragment entry_point_twoOutput entry_point_two( struct entry_point_threeOutput { metal::float4 member_2 [[color(0)]]; }; +struct entry_point_threeArgumentBuffer { + metal::texture2d t1_ [[id(0)]]; + metal::texture2d t2_ [[id(1)]]; + metal::sampler s1_ [[id(2)]]; + metal::sampler s2_ [[id(3)]]; + constant metal::float2& uniformOne [[id(4)]]; + constant metal::float2& uniformTwo [[id(5)]]; +}; fragment entry_point_threeOutput entry_point_three( - metal::texture2d t1_ [[texture(0)]] -, metal::texture2d t2_ [[texture(1)]] -, metal::sampler s2_ [[sampler(1)]] -, constant metal::float2& uniformOne [[buffer(0)]] -, constant metal::float2& uniformTwo [[buffer(1)]] + device entry_point_threeArgumentBuffer& argumentBuffer [[buffer(0)]] ) { constexpr metal::sampler s1_( metal::s_address::clamp_to_edge, @@ -65,10 +76,10 @@ fragment entry_point_threeOutput entry_point_three( metal::min_filter::linear, metal::coord::normalized ); - metal::float2 _e3 = uniformTwo; - metal::float2 _e5 = uniformOne; - metal::float4 _e7 = t1_.sample(s1_, _e3 + _e5); - metal::float2 _e11 = uniformOne; - metal::float4 _e12 = t2_.sample(s2_, _e11); + metal::float2 _e3 = argumentBuffer.uniformTwo; + metal::float2 _e5 = argumentBuffer.uniformOne; + metal::float4 _e7 = argumentBuffer.t1_.sample(argumentBuffer.s1_, _e3 + _e5); + metal::float2 _e11 = argumentBuffer.uniformOne; + metal::float4 _e12 = argumentBuffer.t2_.sample(argumentBuffer.s2_, _e11); return entry_point_threeOutput { _e7 + _e12 }; } diff --git a/tests/out/msl/shadow.msl b/tests/out/msl/shadow.msl index 53f320344a..82a7075569 100644 --- a/tests/out/msl/shadow.msl +++ b/tests/out/msl/shadow.msl @@ -58,20 +58,23 @@ struct vs_mainOutput { metal::float3 world_normal [[user(loc0), center_perspective]]; metal::float4 world_position [[user(loc1), center_perspective]]; }; +struct vs_mainArgumentBuffer { + constant Globals& u_globals [[id(0)]]; + constant Entity& u_entity [[id(1)]]; +}; vertex vs_mainOutput vs_main( vs_mainInput varyings [[stage_in]] -, constant Globals& u_globals [[user(fake0)]] -, constant Entity& u_entity [[user(fake0)]] +, device vs_mainArgumentBuffer& argumentBuffer [[buffer(0)]] ) { const auto position = varyings.position; const auto normal = varyings.normal; VertexOutput out = {}; - metal::float4x4 w = u_entity.world; - metal::float4x4 _e7 = u_entity.world; + metal::float4x4 w = argumentBuffer.u_entity.world; + metal::float4x4 _e7 = argumentBuffer.u_entity.world; metal::float4 world_pos = _e7 * static_cast(position); out.world_normal = metal::float3x3(w[0].xyz, w[1].xyz, w[2].xyz) * static_cast(normal.xyz); out.world_position = world_pos; - metal::float4x4 _e26 = u_globals.view_proj; + metal::float4x4 _e26 = argumentBuffer.u_globals.view_proj; out.proj_position = _e26 * world_pos; VertexOutput _e28 = out; const auto _tmp = _e28; @@ -86,14 +89,17 @@ struct fs_mainInput { struct fs_mainOutput { metal::float4 member_1 [[color(0)]]; }; +struct fs_mainArgumentBuffer { + constant Globals& u_globals [[id(0)]]; + constant Entity& u_entity [[id(1)]]; + device type_6 const& s_lights [[id(2)]]; + metal::depth2d_array t_shadow [[id(3)]]; + metal::sampler sampler_shadow [[id(4)]]; +}; fragment fs_mainOutput fs_main( fs_mainInput varyings_1 [[stage_in]] , metal::float4 proj_position [[position]] -, constant Globals& u_globals [[user(fake0)]] -, constant Entity& u_entity [[user(fake0)]] -, device type_6 const& s_lights [[user(fake0)]] -, metal::depth2d_array t_shadow [[user(fake0)]] -, metal::sampler sampler_shadow [[user(fake0)]] +, device fs_mainArgumentBuffer& argumentBuffer [[buffer(0)]] , constant _mslBufferSizes& _buffer_sizes [[user(fake0)]] ) { const VertexOutput in = { proj_position, varyings_1.world_normal, varyings_1.world_position }; @@ -110,16 +116,16 @@ fragment fs_mainOutput fs_main( } loop_init = false; uint _e7 = i; - uint _e11 = u_globals.num_lights.x; + uint _e11 = argumentBuffer.u_globals.num_lights.x; if (_e7 < metal::min(_e11, c_max_lights)) { } else { break; } { uint _e16 = i; - Light light = s_lights[_e16]; + Light light = argumentBuffer.s_lights[_e16]; uint _e19 = i; - float _e23 = fetch_shadow(_e19, light.proj * in.world_position, t_shadow, sampler_shadow); + float _e23 = fetch_shadow(_e19, light.proj * in.world_position, argumentBuffer.t_shadow, argumentBuffer.sampler_shadow); metal::float3 light_dir = metal::normalize(light.pos.xyz - in.world_position.xyz); float diffuse = metal::max(0.0, metal::dot(normal_1, light_dir)); metal::float3 _e37 = color; @@ -127,7 +133,7 @@ fragment fs_mainOutput fs_main( } } metal::float3 _e42 = color; - metal::float4 _e47 = u_entity.color; + metal::float4 _e47 = argumentBuffer.u_entity.color; return fs_mainOutput { metal::float4(_e42, 1.0) * _e47 }; } @@ -139,14 +145,17 @@ struct fs_main_without_storageInput { struct fs_main_without_storageOutput { metal::float4 member_2 [[color(0)]]; }; +struct fs_main_without_storageArgumentBuffer { + constant Globals& u_globals [[id(0)]]; + constant Entity& u_entity [[id(1)]]; + constant type_7& u_lights [[id(2)]]; + metal::depth2d_array t_shadow [[id(12)]]; + metal::sampler sampler_shadow [[id(13)]]; +}; fragment fs_main_without_storageOutput fs_main_without_storage( fs_main_without_storageInput varyings_2 [[stage_in]] , metal::float4 proj_position_1 [[position]] -, constant Globals& u_globals [[user(fake0)]] -, constant Entity& u_entity [[user(fake0)]] -, constant type_7& u_lights [[user(fake0)]] -, metal::depth2d_array t_shadow [[user(fake0)]] -, metal::sampler sampler_shadow [[user(fake0)]] +, device fs_main_without_storageArgumentBuffer& argumentBuffer [[buffer(0)]] ) { const VertexOutput in_1 = { proj_position_1, varyings_2.world_normal, varyings_2.world_position }; metal::float3 color_1 = {}; @@ -162,16 +171,16 @@ fragment fs_main_without_storageOutput fs_main_without_storage( } loop_init_1 = false; uint _e7 = i_1; - uint _e11 = u_globals.num_lights.x; + uint _e11 = argumentBuffer.u_globals.num_lights.x; if (_e7 < metal::min(_e11, c_max_lights)) { } else { break; } { uint _e16 = i_1; - Light light_1 = u_lights.inner[_e16]; + Light light_1 = argumentBuffer.u_lights.inner[_e16]; uint _e19 = i_1; - float _e23 = fetch_shadow(_e19, light_1.proj * in_1.world_position, t_shadow, sampler_shadow); + float _e23 = fetch_shadow(_e19, light_1.proj * in_1.world_position, argumentBuffer.t_shadow, argumentBuffer.sampler_shadow); metal::float3 light_dir_1 = metal::normalize(light_1.pos.xyz - in_1.world_position.xyz); float diffuse_1 = metal::max(0.0, metal::dot(normal_2, light_dir_1)); metal::float3 _e37 = color_1; @@ -179,6 +188,6 @@ fragment fs_main_without_storageOutput fs_main_without_storage( } } metal::float3 _e42 = color_1; - metal::float4 _e47 = u_entity.color; + metal::float4 _e47 = argumentBuffer.u_entity.color; return fs_main_without_storageOutput { metal::float4(_e42, 1.0) * _e47 }; } diff --git a/tests/out/msl/skybox.msl b/tests/out/msl/skybox.msl index 7b10ea23e7..bab40bdbbf 100644 --- a/tests/out/msl/skybox.msl +++ b/tests/out/msl/skybox.msl @@ -19,9 +19,12 @@ struct vs_mainOutput { metal::float4 position [[position]]; metal::float3 uv [[user(loc0), center_perspective]]; }; +struct vs_mainArgumentBuffer { + constant Data& r_data [[id(0)]]; +}; vertex vs_mainOutput vs_main( uint vertex_index [[vertex_id]] -, constant Data& r_data [[buffer(0)]] +, device vs_mainArgumentBuffer& argumentBuffer [[buffer(0)]] ) { int tmp1_ = {}; int tmp2_ = {}; @@ -30,11 +33,11 @@ vertex vs_mainOutput vs_main( int _e9 = tmp1_; int _e15 = tmp2_; metal::float4 pos = metal::float4((static_cast(_e9) * 4.0) - 1.0, (static_cast(_e15) * 4.0) - 1.0, 0.0, 1.0); - metal::float4 _e27 = r_data.view[0]; - metal::float4 _e32 = r_data.view[1]; - metal::float4 _e37 = r_data.view[2]; + metal::float4 _e27 = argumentBuffer.r_data.view[0]; + metal::float4 _e32 = argumentBuffer.r_data.view[1]; + metal::float4 _e37 = argumentBuffer.r_data.view[2]; metal::float3x3 inv_model_view = metal::transpose(metal::float3x3(_e27.xyz, _e32.xyz, _e37.xyz)); - metal::float4x4 _e43 = r_data.proj_inv; + metal::float4x4 _e43 = argumentBuffer.r_data.proj_inv; metal::float4 unprojected = _e43 * pos; const auto _tmp = VertexOutput {pos, inv_model_view * unprojected.xyz}; return vs_mainOutput { _tmp.position, _tmp.uv }; @@ -47,10 +50,14 @@ struct fs_mainInput { struct fs_mainOutput { metal::float4 member_1 [[color(0)]]; }; +struct fs_mainArgumentBuffer { + metal::texturecube r_texture [[id(0)]]; + metal::sampler r_sampler [[id(1)]]; +}; fragment fs_mainOutput fs_main( fs_mainInput varyings_1 [[stage_in]] , metal::float4 position [[position]] -, metal::texturecube r_texture [[texture(0)]] +, device fs_mainArgumentBuffer& argumentBuffer [[buffer(0)]] ) { constexpr metal::sampler r_sampler( metal::s_address::clamp_to_edge, @@ -61,6 +68,6 @@ fragment fs_mainOutput fs_main( metal::coord::normalized ); const VertexOutput in = { position, varyings_1.uv }; - metal::float4 _e4 = r_texture.sample(r_sampler, in.uv); + metal::float4 _e4 = argumentBuffer.r_texture.sample(argumentBuffer.r_sampler, in.uv); return fs_mainOutput { _e4 }; } diff --git a/tests/out/msl/texture-arg.msl b/tests/out/msl/texture-arg.msl index 5fb9b25649..a492807754 100644 --- a/tests/out/msl/texture-arg.msl +++ b/tests/out/msl/texture-arg.msl @@ -16,10 +16,13 @@ metal::float4 test( struct main_Output { metal::float4 member [[color(0)]]; }; +struct main_ArgumentBuffer { + metal::texture2d Texture [[id(0)]]; + metal::sampler Sampler [[id(1)]]; +}; fragment main_Output main_( - metal::texture2d Texture [[user(fake0)]] -, metal::sampler Sampler [[user(fake0)]] + device main_ArgumentBuffer& argumentBuffer [[buffer(0)]] ) { - metal::float4 _e2 = test(Texture, Sampler); + metal::float4 _e2 = test(argumentBuffer.Texture, argumentBuffer.Sampler); return main_Output { _e2 }; } diff --git a/tests/out/msl/workgroup-var-init.msl b/tests/out/msl/workgroup-var-init.msl index ac300d4337..c34fa859e8 100644 --- a/tests/out/msl/workgroup-var-init.msl +++ b/tests/out/msl/workgroup-var-init.msl @@ -19,10 +19,13 @@ struct WStruct { type_4 atom_arr; }; +struct main_ArgumentBuffer { + device type_1& output [[id(0)]]; +}; kernel void main_( metal::uint3 __local_invocation_id [[thread_position_in_threadgroup]] +, device main_ArgumentBuffer& argumentBuffer [[buffer(0)]] , threadgroup WStruct& w_mem -, device type_1& output [[buffer(0)]] ) { if (metal::all(__local_invocation_id == metal::uint3(0u))) { w_mem.arr = {}; @@ -35,6 +38,6 @@ kernel void main_( } metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); type_1 _e3 = w_mem.arr; - output = _e3; + argumentBuffer.output = _e3; return; } diff --git a/tests/snapshots.rs b/tests/snapshots.rs index a2e03679b2..8bb7d7885f 100644 --- a/tests/snapshots.rs +++ b/tests/snapshots.rs @@ -317,7 +317,7 @@ fn write_output_msl( let (string, tr_info) = msl::write_string(module, info, &options, pipeline_options) .unwrap_or_else(|err| panic!("Metal write failed: {err}")); - for (ep, result) in module.entry_points.iter().zip(tr_info.entry_point_names) { + for (ep, result) in module.entry_points.iter().zip(tr_info.entry_point_info) { if let Err(error) = result { panic!("Failed to translate '{}': {}", ep.name, error); }