From 2774dcb40359349753a8825eca947508e770aa1f Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Thu, 8 Apr 2021 11:00:11 -0400 Subject: [PATCH] [msl] inline some of the types --- src/back/msl/writer.rs | 444 ++++++++++++++++++++++------------ tests/out/boids.msl.snap | 46 ++-- tests/out/collatz.msl.snap | 17 +- tests/out/empty.msl.snap | 1 + tests/out/image-copy.msl.snap | 12 +- tests/out/quad-vert.msl.snap | 46 ++-- tests/out/quad.msl.snap | 27 +-- tests/out/shadow.msl.snap | 46 ++-- tests/out/skybox.msl.snap | 38 ++- 9 files changed, 390 insertions(+), 287 deletions(-) diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index bea412cda3..0c4979a891 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -3,7 +3,7 @@ use super::{ TranslationInfo, }; use crate::{ - arena::Handle, + arena::{Arena, Handle}, proc::{EntryPointIndex, NameKey, Namer, TypeResolution}, valid::{FunctionInfo, GlobalUse, ModuleInfo}, FastHashMap, @@ -31,6 +31,154 @@ impl Display for Level { } } +struct TypeContext<'a> { + handle: Handle, + arena: &'a Arena, + names: &'a FastHashMap, + usage: GlobalUse, + access: crate::StorageAccess, + first_time: bool, +} + +impl<'a> Display for TypeContext<'a> { + fn fmt(&self, out: &mut Formatter<'_>) -> Result<(), FmtError> { + let ty = &self.arena[self.handle]; + if ty.needs_alias() && !self.first_time { + let name = &self.names[&NameKey::Type(self.handle)]; + return write!(out, "{}", name); + } + + match ty.inner { + // work around Metal toolchain bug with `uint` typedef + crate::TypeInner::Scalar { + kind: crate::ScalarKind::Uint, + .. + } => { + write!(out, "metal::uint") + } + crate::TypeInner::Scalar { kind, .. } => { + write!(out, "{}", scalar_kind_string(kind)) + } + crate::TypeInner::Vector { size, kind, .. } => { + write!( + out, + "{}::{}{}", + NAMESPACE, + scalar_kind_string(kind), + vector_size_string(size), + ) + } + crate::TypeInner::Matrix { columns, rows, .. } => { + write!( + out, + "{}::{}{}x{}", + NAMESPACE, + scalar_kind_string(crate::ScalarKind::Float), + vector_size_string(columns), + vector_size_string(rows), + ) + } + crate::TypeInner::Pointer { base, class } => { + let sub = Self { + arena: self.arena, + names: self.names, + handle: base, + usage: self.usage, + access: self.access, + first_time: false, + }; + let class_name = match class.get_name(self.usage) { + Some(name) => name, + None => return Ok(()), + }; + write!(out, "{} {}*", class_name, sub) + } + crate::TypeInner::ValuePointer { + size: None, + kind, + width: _, + class, + } => { + let class_name = match class.get_name(self.usage) { + Some(name) => name, + None => return Ok(()), + }; + write!(out, "{} {}*", class_name, scalar_kind_string(kind),) + } + crate::TypeInner::ValuePointer { + size: Some(size), + kind, + width: _, + class, + } => { + let class_name = match class.get_name(self.usage) { + Some(name) => name, + None => return Ok(()), + }; + write!( + out, + "{} {}::{}{}*", + class_name, + NAMESPACE, + scalar_kind_string(kind), + vector_size_string(size), + ) + } + crate::TypeInner::Array { .. } | crate::TypeInner::Struct { .. } => unreachable!(), + crate::TypeInner::Image { + dim, + arrayed, + class, + } => { + let dim_str = match dim { + crate::ImageDimension::D1 => "1d", + crate::ImageDimension::D2 => "2d", + crate::ImageDimension::D3 => "3d", + crate::ImageDimension::Cube => "cube", + }; + let (texture_str, msaa_str, kind, access) = match class { + crate::ImageClass::Sampled { kind, multi } => { + ("texture", if multi { "_ms" } else { "" }, kind, "sample") + } + crate::ImageClass::Depth => ("depth", "", crate::ScalarKind::Float, "sample"), + crate::ImageClass::Storage(format) => { + let access = if self + .access + .contains(crate::StorageAccess::LOAD | crate::StorageAccess::STORE) + { + "read_write" + } else if self.access.contains(crate::StorageAccess::STORE) { + "write" + } else if self.access.contains(crate::StorageAccess::LOAD) { + "read" + } else { + unreachable!("module is not valid") + }; + ("texture", "", format.into(), access) + } + }; + let base_name = scalar_kind_string(kind); + let array_str = if arrayed { "_array" } else { "" }; + write!( + out, + "{}::{}{}{}{}<{}, {}::access::{}>", + NAMESPACE, + texture_str, + dim_str, + msaa_str, + array_str, + base_name, + NAMESPACE, + access, + ) + } + crate::TypeInner::Sampler { comparison: _ } => { + write!(out, "{}::sampler", NAMESPACE) + } + } + } +} + struct TypedGlobalVariable<'a> { module: &'a crate::Module, names: &'a FastHashMap, @@ -43,7 +191,14 @@ impl<'a> TypedGlobalVariable<'a> { fn try_fmt(&self, out: &mut W) -> Result<(), Error> { let var = &self.module.global_variables[self.handle]; let name = &self.names[&NameKey::GlobalVariable(self.handle)]; - let ty_name = &self.names[&NameKey::Type(var.ty)]; + let ty_name = TypeContext { + handle: var.ty, + arena: &self.module.types, + names: self.names, + usage: self.usage, + access: var.storage_access, + first_time: false, + }; let (space, access, reference) = match var.class.get_name(self.usage) { Some(space) if self.reference => { @@ -143,6 +298,25 @@ impl crate::StorageClass { } } +impl crate::Type { + // Returns `true` if we need to emit a type alias for this type. + fn needs_alias(&self) -> bool { + use crate::TypeInner as Ti; + match self.inner { + // value types are concise enough, we only alias them if they are named + Ti::Scalar { .. } + | Ti::Vector { .. } + | Ti::Matrix { .. } + | Ti::Pointer { .. } + | Ti::ValuePointer { .. } => self.name.is_some(), + // composite types are better to be aliased, regardless of the name + Ti::Struct { .. } | Ti::Array { .. } => true, + // handle types may be different, depending on the global var access, so we always inline them + Ti::Image { .. } | Ti::Sampler { .. } => false, + } + } +} + enum FunctionOrigin { Handle(Handle), EntryPoint(EntryPointIndex), @@ -830,7 +1004,14 @@ impl Writer { ) -> Result<(), Error> { match context.info[handle].ty { TypeResolution::Handle(ty_handle) => { - let ty_name = &self.names[&NameKey::Type(ty_handle)]; + let ty_name = TypeContext { + handle: ty_handle, + arena: &context.module.types, + names: &self.names, + usage: GlobalUse::all(), + access: crate::StorageAccess::empty(), + first_time: false, + }; write!(self.out, "{}", ty_name)?; } TypeResolution::Value(crate::TypeInner::Scalar { kind, .. }) => { @@ -1096,92 +1277,25 @@ impl Writer { fn write_type_defs(&mut self, module: &crate::Module) -> Result<(), Error> { for (handle, ty) in module.types.iter() { + if !ty.needs_alias() { + continue; + } let name = &self.names[&NameKey::Type(handle)]; let global_use = GlobalUse::all(); //TODO match ty.inner { - // work around Metal toolchain bug with `uint` typedef - crate::TypeInner::Scalar { - kind: crate::ScalarKind::Uint, - .. - } => { - writeln!(self.out, "typedef metal::uint {};", name)?; - } - crate::TypeInner::Scalar { kind, .. } => { - writeln!(self.out, "typedef {} {};", scalar_kind_string(kind), name)?; - } - crate::TypeInner::Vector { size, kind, .. } => { - writeln!( - self.out, - "typedef {}::{}{} {};", - NAMESPACE, - scalar_kind_string(kind), - vector_size_string(size), - name - )?; - } - crate::TypeInner::Matrix { columns, rows, .. } => { - writeln!( - self.out, - "typedef {}::{}{}x{} {};", - NAMESPACE, - scalar_kind_string(crate::ScalarKind::Float), - vector_size_string(columns), - vector_size_string(rows), - name - )?; - } - crate::TypeInner::Pointer { base, class } => { - let base_name = &self.names[&NameKey::Type(base)]; - let class_name = match class.get_name(global_use) { - Some(name) => name, - None => continue, - }; - writeln!(self.out, "typedef {} {} *{};", class_name, base_name, name)?; - } - crate::TypeInner::ValuePointer { - size: None, - kind, - width: _, - class, - } => { - let class_name = match class.get_name(global_use) { - Some(name) => name, - None => continue, - }; - writeln!( - self.out, - "typedef {} {} *{};", - class_name, - scalar_kind_string(kind), - name - )?; - } - crate::TypeInner::ValuePointer { - size: Some(size), - kind, - width: _, - class, - } => { - let class_name = match class.get_name(global_use) { - Some(name) => name, - None => continue, - }; - writeln!( - self.out, - "typedef {} {}::{}{} {};", - class_name, - NAMESPACE, - scalar_kind_string(kind), - vector_size_string(size), - name - )?; - } crate::TypeInner::Array { base, size, stride: _, } => { - let base_name = &self.names[&NameKey::Type(base)]; + let base_name = TypeContext { + handle: base, + arena: &module.types, + names: &self.names, + usage: global_use, + access: crate::StorageAccess::empty(), + first_time: false, + }; let size_str = match size { crate::ArraySize::Constant(const_handle) => { &self.names[&NameKey::Constant(const_handle)] @@ -1197,68 +1311,28 @@ impl Writer { writeln!(self.out, "struct {} {{", name)?; for (index, member) in members.iter().enumerate() { let member_name = &self.names[&NameKey::StructMember(handle, index as u32)]; - let base_name = &self.names[&NameKey::Type(member.ty)]; + let base_name = TypeContext { + handle: member.ty, + arena: &module.types, + names: &self.names, + usage: global_use, + access: crate::StorageAccess::empty(), + first_time: false, + }; writeln!(self.out, "{}{} {};", INDENT, base_name, member_name)?; } writeln!(self.out, "}};")?; } - crate::TypeInner::Image { - dim, - arrayed, - class, - } => { - let dim_str = match dim { - crate::ImageDimension::D1 => "1d", - crate::ImageDimension::D2 => "2d", - crate::ImageDimension::D3 => "3d", - crate::ImageDimension::Cube => "cube", + _ => { + let ty_name = TypeContext { + handle, + arena: &module.types, + names: &self.names, + usage: global_use, + access: crate::StorageAccess::empty(), + first_time: true, }; - let (texture_str, msaa_str, kind, access) = match class { - crate::ImageClass::Sampled { kind, multi } => { - ("texture", if multi { "_ms" } else { "" }, kind, "sample") - } - crate::ImageClass::Depth => { - ("depth", "", crate::ScalarKind::Float, "sample") - } - crate::ImageClass::Storage(format) => { - let (_, global) = module - .global_variables - .iter() - .find(|&(_, ref var)| var.ty == handle) - .expect("Unable to find a global variable using the image type"); - let access = if global - .storage_access - .contains(crate::StorageAccess::LOAD | crate::StorageAccess::STORE) - { - "read_write" - } else if global.storage_access.contains(crate::StorageAccess::STORE) { - "write" - } else if global.storage_access.contains(crate::StorageAccess::LOAD) { - "read" - } else { - return Err(Error::Validation); - }; - ("texture", "", format.into(), access) - } - }; - let base_name = scalar_kind_string(kind); - let array_str = if arrayed { "_array" } else { "" }; - writeln!( - self.out, - "typedef {}::{}{}{}{}<{}, {}::access::{}> {};", - NAMESPACE, - texture_str, - dim_str, - msaa_str, - array_str, - base_name, - NAMESPACE, - access, - name - )?; - } - crate::TypeInner::Sampler { comparison: _ } => { - writeln!(self.out, "typedef {}::sampler {};", NAMESPACE, name)?; + writeln!(self.out, "typedef {} {};", ty_name, name)?; } } } @@ -1305,7 +1379,14 @@ impl Writer { crate::ConstantInner::Scalar { .. } => {} crate::ConstantInner::Composite { ty, ref components } => { let name = &self.names[&NameKey::Constant(handle)]; - let ty_name = &self.names[&NameKey::Type(ty)]; + let ty_name = TypeContext { + handle: ty, + arena: &module.types, + names: &self.names, + usage: GlobalUse::empty(), + access: crate::StorageAccess::empty(), + first_time: false, + }; write!(self.out, "constexpr constant {} {} = {{", ty_name, name,)?; for (i, &sub_handle) in components.iter().enumerate() { let separator = if i != 0 { ", " } else { "" }; @@ -1415,16 +1496,36 @@ impl Writer { } } + writeln!(self.out)?; let fun_name = &self.names[&NameKey::Function(fun_handle)]; - let result_type_name = match fun.result { - Some(ref result) => &self.names[&NameKey::Type(result.ty)], - None => "void", - }; - writeln!(self.out, "{} {}(", result_type_name, fun_name)?; + match fun.result { + Some(ref result) => { + let ty_name = TypeContext { + handle: result.ty, + arena: &module.types, + names: &self.names, + usage: GlobalUse::empty(), + access: crate::StorageAccess::empty(), + first_time: false, + }; + write!(self.out, "{}", ty_name)?; + } + None => { + write!(self.out, "void")?; + } + } + writeln!(self.out, " {}(", fun_name)?; for (index, arg) in fun.arguments.iter().enumerate() { let name = &self.names[&NameKey::FunctionArgument(fun_handle, index as u32)]; - let param_type_name = &self.names[&NameKey::Type(arg.ty)]; + let param_type_name = TypeContext { + handle: arg.ty, + arena: &module.types, + names: &self.names, + usage: GlobalUse::empty(), + access: crate::StorageAccess::empty(), + first_time: false, + }; let separator = separate(!pass_through_globals.is_empty() || index + 1 != fun.arguments.len()); writeln!( @@ -1449,7 +1550,14 @@ impl Writer { writeln!(self.out, ") {{")?; for (local_handle, local) in fun.local_variables.iter() { - let ty_name = &self.names[&NameKey::Type(local.ty)]; + let ty_name = TypeContext { + handle: local.ty, + arena: &module.types, + names: &self.names, + usage: GlobalUse::empty(), + access: crate::StorageAccess::empty(), + first_time: false, + }; let local_name = &self.names[&NameKey::FunctionLocal(fun_handle, local_handle)]; write!(self.out, "{}{} {}", INDENT, ty_name, local_name)?; if let Some(value) = local.init { @@ -1473,7 +1581,6 @@ impl Writer { self.named_expressions.clear(); self.put_block(Level(1), &fun.body, &context)?; writeln!(self.out, "}}")?; - writeln!(self.out)?; } let mut info = TranslationInfo { @@ -1503,6 +1610,7 @@ impl Writer { } } + writeln!(self.out)?; let fun_name = &self.names[&NameKey::EntryPoint(ep_index as _)]; info.entry_point_names.push(Ok(fun_name.clone())); @@ -1558,9 +1666,16 @@ impl Writer { }; varying_count += 1; let name = &self.names[&name_key]; - let type_name = &self.names[&NameKey::Type(ty)]; + let ty_name = TypeContext { + handle: ty, + arena: &module.types, + names: &self.names, + usage: GlobalUse::empty(), + access: crate::StorageAccess::empty(), + first_time: false, + }; let resolved = options.resolve_local_binding(binding, in_mode)?; - write!(self.out, "{}{} {}", INDENT, type_name, name)?; + write!(self.out, "{}{} {}", INDENT, ty_name, name)?; resolved.try_fmt_decorated(&mut self.out, "")?; writeln!(self.out, ";")?; } @@ -1592,7 +1707,14 @@ impl Writer { } writeln!(self.out, "struct {} {{", stage_out_name)?; for (name, ty, binding) in result_members { - let type_name = &self.names[&NameKey::Type(ty)]; + let ty_name = TypeContext { + handle: ty, + arena: &module.types, + names: &self.names, + usage: GlobalUse::empty(), + access: crate::StorageAccess::empty(), + first_time: false, + }; let binding = binding.ok_or(Error::Validation)?; if !pipeline_options.allow_point_size && *binding == crate::Binding::BuiltIn(crate::BuiltIn::PointSize) @@ -1600,7 +1722,7 @@ impl Writer { continue; } let resolved = options.resolve_local_binding(binding, out_mode)?; - write!(self.out, "{}{} {}", INDENT, type_name, name)?; + write!(self.out, "{}{} {}", INDENT, ty_name, name)?; resolved.try_fmt_decorated(&mut self.out, "")?; writeln!(self.out, ";")?; } @@ -1626,7 +1748,14 @@ impl Writer { _ => continue, }; let name = &self.names[&name_key]; - let type_name = &self.names[&NameKey::Type(ty)]; + let ty_name = TypeContext { + handle: ty, + arena: &module.types, + names: &self.names, + usage: GlobalUse::empty(), + access: crate::StorageAccess::empty(), + first_time: false, + }; let resolved = options.resolve_local_binding(binding, in_mode)?; let separator = if is_first_argument { is_first_argument = false; @@ -1634,7 +1763,7 @@ impl Writer { } else { ',' }; - write!(self.out, "{} {} {}", separator, type_name, name)?; + write!(self.out, "{} {} {}", separator, ty_name, name)?; resolved.try_fmt_decorated(&mut self.out, "\n")?; } for (handle, var) in module.global_variables.iter() { @@ -1763,7 +1892,14 @@ impl Writer { //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 = &self.names[&NameKey::Type(local.ty)]; + let ty_name = TypeContext { + handle: local.ty, + arena: &module.types, + names: &self.names, + usage: GlobalUse::empty(), + access: crate::StorageAccess::empty(), + first_time: false, + }; write!(self.out, "{}{} {}", INDENT, ty_name, name)?; if let Some(value) = local.init { let value_str = &self.names[&NameKey::Constant(value)]; diff --git a/tests/out/boids.msl.snap b/tests/out/boids.msl.snap index e690e2ccdf..9f302992a6 100644 --- a/tests/out/boids.msl.snap +++ b/tests/out/boids.msl.snap @@ -14,46 +14,42 @@ constexpr constant unsigned const_1u = 1u; constexpr constant float const_1f = 1.0; constexpr constant float const_0_10f = 0.1; constexpr constant float const_n1f = -1.0; -typedef metal::uint type; -typedef metal::float2 type1; struct Particle { - type1 pos; - type1 vel; + metal::float2 pos; + metal::float2 vel; }; -typedef float type2; struct SimParams { - type2 deltaT; - type2 rule1Distance; - type2 rule2Distance; - type2 rule3Distance; - type2 rule1Scale; - type2 rule2Scale; - type2 rule3Scale; + float deltaT; + float rule1Distance; + float rule2Distance; + float rule3Distance; + float rule1Scale; + float rule2Scale; + float rule3Scale; }; typedef Particle type3[1]; struct Particles { type3 particles; }; -typedef metal::uint3 type4; -typedef int type5; + struct main1Input { }; kernel void main1( - type4 global_invocation_id [[thread_position_in_grid]] + metal::uint3 global_invocation_id [[thread_position_in_grid]] , constant SimParams& params [[buffer(0)]] , constant Particles& particlesSrc [[buffer(1)]] , device Particles& particlesDst [[buffer(2)]] ) { - type1 vPos; - type1 vVel; - type1 cMass; - type1 cVel; - type1 colVel; - type5 cMassCount = const_0i; - type5 cVelCount = const_0i; - type1 pos1; - type1 vel1; - type i = const_0u; + metal::float2 vPos; + metal::float2 vVel; + metal::float2 cMass; + metal::float2 cVel; + metal::float2 colVel; + int cMassCount = const_0i; + int cVelCount = const_0i; + metal::float2 pos1; + metal::float2 vel1; + metal::uint i = const_0u; if (global_invocation_id.x >= NUM_PARTICLES) { return; } diff --git a/tests/out/collatz.msl.snap b/tests/out/collatz.msl.snap index 1a0e7d95d0..2aa4fe3446 100644 --- a/tests/out/collatz.msl.snap +++ b/tests/out/collatz.msl.snap @@ -9,17 +9,16 @@ constexpr constant unsigned const_0u = 0u; constexpr constant unsigned const_1u = 1u; constexpr constant unsigned const_2u = 2u; constexpr constant unsigned const_3u = 3u; -typedef metal::uint type; -typedef type type1[1]; +typedef metal::uint type1[1]; struct PrimeIndices { type1 data; }; -typedef metal::uint3 type2; -type collatz_iterations( - type n_base + +metal::uint collatz_iterations( + metal::uint n_base ) { - type n; - type i = const_0u; + metal::uint n; + metal::uint i = const_0u; n = n_base; while(true) { if (n <= const_1u) { @@ -38,10 +37,10 @@ type collatz_iterations( struct main1Input { }; kernel void main1( - type2 global_id [[thread_position_in_grid]] + metal::uint3 global_id [[thread_position_in_grid]] , device PrimeIndices& v_indices [[user(fake0)]] ) { - type _expr9 = collatz_iterations(v_indices.data[global_id.x]); + metal::uint _expr9 = collatz_iterations(v_indices.data[global_id.x]); v_indices.data[global_id.x] = _expr9; return; } diff --git a/tests/out/empty.msl.snap b/tests/out/empty.msl.snap index 6cbbbb4bdd..9d081189ae 100644 --- a/tests/out/empty.msl.snap +++ b/tests/out/empty.msl.snap @@ -5,6 +5,7 @@ expression: msl #include #include + kernel void main1( ) { return; diff --git a/tests/out/image-copy.msl.snap b/tests/out/image-copy.msl.snap index 6584e9463e..3b9ca6f741 100644 --- a/tests/out/image-copy.msl.snap +++ b/tests/out/image-copy.msl.snap @@ -7,17 +7,13 @@ expression: msl constexpr constant int const_10i = 10; constexpr constant int const_20i = 20; -typedef metal::texture2d type; -typedef metal::texture1d type1; -typedef metal::uint3 type2; -typedef metal::uint2 type3; -typedef metal::int2 type4; + struct main1Input { }; kernel void main1( - type2 local_id [[thread_position_in_threadgroup]] -, type image_src [[user(fake0)]] -, type1 image_dst [[user(fake0)]] + metal::uint3 local_id [[thread_position_in_threadgroup]] +, metal::texture2d image_src [[user(fake0)]] +, metal::texture1d image_dst [[user(fake0)]] ) { metal::int2 _expr12 = (int2(image_src.get_width(), image_src.get_height()) * static_cast(metal::uint2(local_id.x, local_id.y))) % metal::int2(const_10i, const_20i); metal::uint4 _expr13 = image_src.read(metal::uint2(_expr12)); diff --git a/tests/out/quad-vert.msl.snap b/tests/out/quad-vert.msl.snap index 16b3529771..a871820759 100644 --- a/tests/out/quad-vert.msl.snap +++ b/tests/out/quad-vert.msl.snap @@ -13,57 +13,49 @@ constexpr constant unsigned const_1u = 1u; constexpr constant int const_0i1 = 0; constexpr constant float const_0f = 0.0; constexpr constant float const_1f = 1.0; -typedef float type; -typedef metal::float2 type1; -typedef thread type1 *type2; -typedef thread type1 *type3; -typedef metal::float4 type4; -typedef metal::uint type5; -typedef type type6[const_1u]; +typedef float type6[const_1u]; struct gl_PerVertex { - type4 gl_Position; - type gl_PointSize; + metal::float4 gl_Position; + float gl_PointSize; type6 gl_ClipDistance; type6 gl_CullDistance; }; -typedef thread gl_PerVertex *type7; -typedef int type8; -typedef thread type4 *type9; struct type10 { - type1 member; - type4 gl_Position1; - type gl_PointSize1; + metal::float2 member; + metal::float4 gl_Position1; + float gl_PointSize1; type6 gl_ClipDistance1; }; + void main1( - thread type1& v_uv, - thread type1 const& a_uv, + thread metal::float2& v_uv, + thread metal::float2 const& a_uv, thread gl_PerVertex& _, - thread type1 const& a_pos + thread metal::float2 const& a_pos ) { v_uv = a_uv; - type1 _expr13 = a_pos; + metal::float2 _expr13 = a_pos; _.gl_Position = metal::float4(_expr13.x, _expr13.y, const_0f, const_1f); return; } struct main2Input { - type1 a_uv1 [[attribute(1)]]; - type1 a_pos1 [[attribute(0)]]; + metal::float2 a_uv1 [[attribute(1)]]; + metal::float2 a_pos1 [[attribute(0)]]; }; struct main2Output { - type1 member [[user(loc0)]]; - type4 gl_Position1 [[position]]; - type gl_PointSize1 [[point_size]]; + metal::float2 member [[user(loc0)]]; + metal::float4 gl_Position1 [[position]]; + float gl_PointSize1 [[point_size]]; type6 gl_ClipDistance1 [[clip_distance]]; }; vertex main2Output main2( main2Input varyings [[stage_in]] ) { - type1 v_uv = {}; - type1 a_uv = {}; + metal::float2 v_uv = {}; + metal::float2 a_uv = {}; gl_PerVertex _ = {}; - type1 a_pos = {}; + metal::float2 a_pos = {}; const auto a_uv1 = varyings.a_uv1; const auto a_pos1 = varyings.a_pos1; a_uv = a_uv1; diff --git a/tests/out/quad.msl.snap b/tests/out/quad.msl.snap index 47a48a2d18..434a62aa64 100644 --- a/tests/out/quad.msl.snap +++ b/tests/out/quad.msl.snap @@ -8,22 +8,18 @@ expression: msl constexpr constant float c_scale = 1.2; constexpr constant float const_0f = 0.0; constexpr constant float const_1f = 1.0; -typedef float type; -typedef metal::float2 type1; -typedef metal::float4 type2; struct VertexOutput { - type1 uv; - type2 position; + metal::float2 uv; + metal::float4 position; }; -typedef metal::texture2d type3; -typedef metal::sampler type4; + struct main1Input { - type1 pos [[attribute(0)]]; - type1 uv1 [[attribute(1)]]; + metal::float2 pos [[attribute(0)]]; + metal::float2 uv1 [[attribute(1)]]; }; struct main1Output { - type1 uv [[user(loc0)]]; - type2 position [[position]]; + metal::float2 uv [[user(loc0)]]; + metal::float4 position [[position]]; }; vertex main1Output main1( main1Input varyings [[stage_in]] @@ -37,16 +33,17 @@ vertex main1Output main1( return main1Output { _tmp.uv, _tmp.position }; } + struct main2Input { - type1 uv2 [[user(loc0)]]; + metal::float2 uv2 [[user(loc0)]]; }; struct main2Output { - type2 member1 [[color(0)]]; + metal::float4 member1 [[color(0)]]; }; fragment main2Output main2( main2Input varyings1 [[stage_in]] -, type3 u_texture [[user(fake0)]] -, type4 u_sampler [[user(fake0)]] +, metal::texture2d u_texture [[user(fake0)]] +, metal::sampler u_sampler [[user(fake0)]] ) { const auto uv2 = varyings1.uv2; metal::float4 _expr4 = u_texture.sample(u_sampler, uv2); diff --git a/tests/out/shadow.msl.snap b/tests/out/shadow.msl.snap index 94cd7982e8..057daff8e8 100644 --- a/tests/out/shadow.msl.snap +++ b/tests/out/shadow.msl.snap @@ -13,33 +13,25 @@ constexpr constant float const_0_05f = 0.05; constexpr constant unsigned c_max_lights = 10u; constexpr constant unsigned const_0u = 0u; constexpr constant unsigned const_1u = 1u; -typedef metal::uint4 type; struct Globals { - type num_lights; + metal::uint4 num_lights; }; -typedef metal::float4x4 type1; -typedef metal::float4 type2; struct Light { - type1 proj; - type2 pos; - type2 color; + metal::float4x4 proj; + metal::float4 pos; + metal::float4 color; }; typedef Light type3[1]; struct Lights { type3 data; }; -typedef metal::depth2d_array type4; -typedef metal::sampler type5; -typedef metal::uint type6; -typedef float type7; -typedef metal::float2 type8; -typedef metal::float3 type9; -constexpr constant type9 c_ambient = {const_0_05f, const_0_05f, const_0_05f}; -type7 fetch_shadow( - type6 light_id, - type2 homogeneous_coords, - type4 t_shadow, - type5 sampler_shadow +constexpr constant metal::float3 c_ambient = {const_0_05f, const_0_05f, const_0_05f}; + +float fetch_shadow( + metal::uint light_id, + metal::float4 homogeneous_coords, + metal::depth2d_array t_shadow, + metal::sampler sampler_shadow ) { if (homogeneous_coords.w <= const_0f) { return const_1f; @@ -50,23 +42,23 @@ type7 fetch_shadow( } struct fs_mainInput { - type9 raw_normal [[user(loc0)]]; - type2 position [[user(loc1)]]; + metal::float3 raw_normal [[user(loc0)]]; + metal::float4 position [[user(loc1)]]; }; struct fs_mainOutput { - type2 member [[color(0)]]; + metal::float4 member [[color(0)]]; }; fragment fs_mainOutput fs_main( fs_mainInput varyings [[stage_in]] , constant Globals& u_globals [[user(fake0)]] , constant Lights& s_lights [[user(fake0)]] -, type4 t_shadow [[user(fake0)]] -, type5 sampler_shadow [[user(fake0)]] +, metal::depth2d_array t_shadow [[user(fake0)]] +, metal::sampler sampler_shadow [[user(fake0)]] ) { const auto raw_normal = varyings.raw_normal; const auto position = varyings.position; - type9 color1 = c_ambient; - type6 i = const_0u; + metal::float3 color1 = c_ambient; + metal::uint i = const_0u; bool loop_init = true; while(true) { if (!loop_init) { @@ -77,7 +69,7 @@ fragment fs_mainOutput fs_main( break; } Light _expr21 = s_lights.data[i]; - type7 _expr25 = fetch_shadow(i, _expr21.proj * position, t_shadow, sampler_shadow); + float _expr25 = fetch_shadow(i, _expr21.proj * position, t_shadow, sampler_shadow); color1 = color1 + ((_expr25 * metal::max(const_0f, metal::dot(metal::normalize(raw_normal), metal::normalize(metal::float3(_expr21.pos.x, _expr21.pos.y, _expr21.pos.z) - metal::float3(position.x, position.y, position.z))))) * metal::float3(_expr21.color.x, _expr21.color.y, _expr21.color.z)); } return fs_mainOutput { metal::float4(color1, const_1f) }; diff --git a/tests/out/skybox.msl.snap b/tests/out/skybox.msl.snap index 5fde024778..96bbabe890 100644 --- a/tests/out/skybox.msl.snap +++ b/tests/out/skybox.msl.snap @@ -10,38 +10,31 @@ constexpr constant int const_1i = 1; constexpr constant float const_4f = 4.0; constexpr constant float const_1f = 1.0; constexpr constant float const_0f = 0.0; -typedef metal::float4 type; -typedef metal::float3 type1; struct VertexOutput { - type position; - type1 uv; + metal::float4 position; + metal::float3 uv; }; -typedef metal::float4x4 type2; struct Data { - type2 proj_inv; - type2 view; + metal::float4x4 proj_inv; + metal::float4x4 view; }; -typedef metal::uint type3; -typedef int type4; -typedef metal::float3x3 type5; -typedef metal::texturecube type6; -typedef metal::sampler type7; + struct vs_mainInput { }; struct vs_mainOutput { - type position [[position]]; - type1 uv [[user(loc0)]]; + metal::float4 position [[position]]; + metal::float3 uv [[user(loc0)]]; }; vertex vs_mainOutput vs_main( - type3 vertex_index [[vertex_id]] + metal::uint vertex_index [[vertex_id]] , constant Data& r_data [[buffer(0)]] ) { - type4 tmp1_; - type4 tmp2_; + int tmp1_; + int tmp2_; VertexOutput out; tmp1_ = static_cast(vertex_index) / const_2i; tmp2_ = static_cast(vertex_index) & const_1i; - type _expr24 = metal::float4((static_cast(tmp1_) * const_4f) - const_1f, (static_cast(tmp2_) * const_4f) - const_1f, const_0f, const_1f); + metal::float4 _expr24 = metal::float4((static_cast(tmp1_) * const_4f) - const_1f, (static_cast(tmp2_) * const_4f) - const_1f, const_0f, const_1f); metal::float4 _expr50 = r_data.proj_inv * _expr24; out.uv = metal::transpose(metal::float3x3(metal::float3(r_data.view[0].x, r_data.view[0].y, r_data.view[0].z), metal::float3(r_data.view[1].x, r_data.view[1].y, r_data.view[1].z), metal::float3(r_data.view[2].x, r_data.view[2].y, r_data.view[2].z))) * metal::float3(_expr50.x, _expr50.y, _expr50.z); out.position = _expr24; @@ -49,16 +42,17 @@ vertex vs_mainOutput vs_main( return vs_mainOutput { _tmp.position, _tmp.uv }; } + struct fs_mainInput { - type1 uv [[user(loc0)]]; + metal::float3 uv [[user(loc0)]]; }; struct fs_mainOutput { - type member1 [[color(0)]]; + metal::float4 member1 [[color(0)]]; }; fragment fs_mainOutput fs_main( fs_mainInput varyings1 [[stage_in]] -, type position [[position]] -, type6 r_texture [[texture(0)]] +, metal::float4 position [[position]] +, metal::texturecube r_texture [[texture(0)]] ) { constexpr metal::sampler r_sampler( metal::s_address::clamp_to_edge,