diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index c2b66fe62d..feaa010986 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -9,6 +9,8 @@ use std::{ io::Write, }; +const NAMESPACE: &str = "metal"; + struct Level(usize); impl Level { fn next(&self) -> Self { @@ -161,9 +163,12 @@ impl Writer { let name = &self.names[&NameKey::StructMember(base_ty, index)]; write!(self.out, ".{}", name)?; } - crate::TypeInner::Matrix { .. } | crate::TypeInner::Vector { .. } => { + crate::TypeInner::Vector { .. } => { write!(self.out, ".{}", COMPONENTS[index as usize])?; } + crate::TypeInner::Matrix { .. } => { + write!(self.out, "[{}]", index)?; + } crate::TypeInner::Array { .. } => { write!(self.out, "[{}]", index)?; } @@ -176,17 +181,30 @@ impl Writer { crate::Expression::Compose { ty, ref components } => { let inner = &context.module.types[ty].inner; match *inner { + crate::TypeInner::Scalar { width: 4, kind } if components.len() == 1 => { + write!(self.out, "{}", scalar_kind_string(kind),)?; + self.put_call("", components, context)?; + } crate::TypeInner::Vector { size, kind, .. } => { write!( self.out, - "{}{}", + "{}::{}{}", + NAMESPACE, scalar_kind_string(kind), vector_size_string(size) )?; self.put_call("", components, context)?; } - crate::TypeInner::Scalar { width: 4, kind } if components.len() == 1 => { - write!(self.out, "{}", scalar_kind_string(kind),)?; + crate::TypeInner::Matrix { columns, rows, .. } => { + let kind = crate::ScalarKind::Float; + write!( + self.out, + "{}::{}{}x{}", + NAMESPACE, + scalar_kind_string(kind), + vector_size_string(columns), + vector_size_string(rows) + )?; self.put_call("", components, context)?; } _ => return Err(Error::UnsupportedCompose(ty)), @@ -666,7 +684,8 @@ impl Writer { crate::TypeInner::Vector { size, kind, .. } => { write!( self.out, - "typedef {}{} {}", + "typedef {}::{}{} {}", + NAMESPACE, scalar_kind_string(kind), vector_size_string(size), name @@ -675,7 +694,8 @@ impl Writer { crate::TypeInner::Matrix { columns, rows, .. } => { write!( self.out, - "typedef {}{}x{} {}", + "typedef {}::{}{}x{} {}", + NAMESPACE, scalar_kind_string(crate::ScalarKind::Float), vector_size_string(columns), vector_size_string(rows), @@ -744,7 +764,7 @@ impl Writer { crate::ImageDimension::D1 => "1d", crate::ImageDimension::D2 => "2d", crate::ImageDimension::D3 => "3d", - crate::ImageDimension::Cube => "Cube", + crate::ImageDimension::Cube => "cube", }; let (texture_str, msaa_str, kind, access) = match class { crate::ImageClass::Sampled { kind, multi } => { @@ -778,12 +798,20 @@ impl Writer { let array_str = if arrayed { "_array" } else { "" }; write!( self.out, - "typedef {}{}{}{}<{}, access::{}> {}", - texture_str, dim_str, msaa_str, array_str, base_name, access, name + "typedef {}::{}{}{}{}<{}, {}::access::{}> {}", + NAMESPACE, + texture_str, + dim_str, + msaa_str, + array_str, + base_name, + NAMESPACE, + access, + name )?; } crate::TypeInner::Sampler { comparison: _ } => { - write!(self.out, "typedef sampler {}", name)?; + write!(self.out, "typedef {}::sampler {}", NAMESPACE, name)?; } } writeln!(self.out, ";")?; diff --git a/tests/snapshots/snapshots__boids.msl.snap b/tests/snapshots/snapshots__boids.msl.snap index 2e3efb40c2..1c1b6800f1 100644 --- a/tests/snapshots/snapshots__boids.msl.snap +++ b/tests/snapshots/snapshots__boids.msl.snap @@ -5,8 +5,8 @@ expression: msl #include #include -typedef float2 type; -typedef float4 type1; +typedef metal::float2 type; +typedef metal::float4 type1; typedef float type2; struct Particle { type pos; @@ -26,7 +26,7 @@ typedef Particle type4[5]; struct Particles { type4 particles; }; -typedef uint3 type5; +typedef metal::uint3 type5; typedef int type6; struct main1Input { @@ -41,7 +41,7 @@ vertex main1Output main1( main1Input input [[stage_in]] ) { main1Output output; - output.gl_Position = float4(float2(input.a_pos.x * metal::cos(-metal::atan2(input.a_particleVel.x, input.a_particleVel.y)) - input.a_pos.y * metal::sin(-metal::atan2(input.a_particleVel.x, input.a_particleVel.y)), input.a_pos.x * metal::sin(-metal::atan2(input.a_particleVel.x, input.a_particleVel.y)) + input.a_pos.y * metal::cos(-metal::atan2(input.a_particleVel.x, input.a_particleVel.y))) + input.a_particlePos, 0.0, 1.0); + output.gl_Position = metal::float4(metal::float2(input.a_pos.x * metal::cos(-metal::atan2(input.a_particleVel.x, input.a_particleVel.y)) - input.a_pos.y * metal::sin(-metal::atan2(input.a_particleVel.x, input.a_particleVel.y)), input.a_pos.x * metal::sin(-metal::atan2(input.a_particleVel.x, input.a_particleVel.y)) + input.a_pos.y * metal::cos(-metal::atan2(input.a_particleVel.x, input.a_particleVel.y))) + input.a_particlePos, 0.0, 1.0); return output; } struct main2Input { @@ -53,7 +53,7 @@ fragment main2Output main2( main2Input input [[stage_in]] ) { main2Output output; - output.fragColor = float4(1.0, 1.0, 1.0, 1.0); + output.fragColor = metal::float4(1.0, 1.0, 1.0, 1.0); return output; } kernel void main3( @@ -76,9 +76,9 @@ kernel void main3( } vPos = particlesA.particles[gl_GlobalInvocationID.x].pos; vVel = particlesA.particles[gl_GlobalInvocationID.x].vel; - cMass = float2(0.0, 0.0); - cVel = float2(0.0, 0.0); - colVel = float2(0.0, 0.0); + cMass = metal::float2(0.0, 0.0); + cVel = metal::float2(0.0, 0.0); + colVel = metal::float2(0.0, 0.0); while(true) { if (i >= static_cast(5)) { break; @@ -86,8 +86,8 @@ kernel void main3( if (i == gl_GlobalInvocationID.x) { continue; } - pos1 = float2(particlesA.particles[i].pos.x, particlesA.particles[i].pos.y); - vel1 = float2(particlesA.particles[i].vel.x, particlesA.particles[i].vel.y); + pos1 = metal::float2(particlesA.particles[i].pos.x, particlesA.particles[i].pos.y); + vel1 = metal::float2(particlesA.particles[i].vel.x, particlesA.particles[i].vel.y); if (metal::distance(pos1, vPos) < params.rule1Distance) { cMass = cMass + pos1; cMassCount = cMassCount + 1; @@ -101,10 +101,10 @@ kernel void main3( } } if (cMassCount == 0) { - cMass = cMass / float2(cMassCount, cMassCount) + vPos; + cMass = cMass / metal::float2(cMassCount, cMassCount) + vPos; } if (cVelCount == 0) { - cVel = cVel / float2(cVelCount, cVelCount); + cVel = cVel / metal::float2(cVelCount, cVelCount); } vVel = vVel + cMass * params.rule1Scale + colVel * params.rule2Scale + cVel * params.rule3Scale; vVel = metal::normalize(vVel) * metal::clamp(metal::length(vVel), 0.0, 0.1); diff --git a/tests/snapshots/snapshots__quad.msl.snap b/tests/snapshots/snapshots__quad.msl.snap index 8fe9dfede3..da25574b10 100644 --- a/tests/snapshots/snapshots__quad.msl.snap +++ b/tests/snapshots/snapshots__quad.msl.snap @@ -6,11 +6,11 @@ expression: msl #include typedef float type; -typedef float2 type1; -typedef float4 type2; -typedef texture2d type3; -typedef sampler type4; -typedef int2 type5; +typedef metal::float2 type1; +typedef metal::float4 type2; +typedef metal::texture2d type3; +typedef metal::sampler type4; +typedef metal::int2 type5; struct main1Input { type1 a_pos [[attribute(0)]]; @@ -25,7 +25,7 @@ vertex main1Output main1( ) { main1Output output; output.v_uv = input.a_uv; - output.o_position = float4(1.2 * input.a_pos, 0.0, 1.0); + output.o_position = metal::float4(1.2 * input.a_pos, 0.0, 1.0); return output; } struct main2Input { diff --git a/tests/snapshots/snapshots__simple.msl.snap b/tests/snapshots/snapshots__simple.msl.snap index 01314fdb2a..0bf6130e99 100644 --- a/tests/snapshots/snapshots__simple.msl.snap +++ b/tests/snapshots/snapshots__simple.msl.snap @@ -5,7 +5,7 @@ expression: msl #include #include -typedef float4 type; +typedef metal::float4 type; typedef int type1; struct main1Input { @@ -17,7 +17,7 @@ vertex main1Output main1( main1Input input [[stage_in]] ) { main1Output output; - output.o_position = float4(1); + output.o_position = metal::float4(1); return output; }