diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index c0d581693e..61bc1d4338 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -183,7 +183,7 @@ impl Writer { if i != 0 { write!(self.out, ", ")?; } - self.put_expression(handle, context)?; + self.put_expression(handle, context, true)?; } write!(self.out, ")")?; Ok(()) @@ -196,10 +196,10 @@ impl Writer { level: Option>, context: &ExpressionContext, ) -> Result<(), Error> { - self.put_expression(image, context)?; + self.put_expression(image, context, false)?; write!(self.out, ".get_{}(", query)?; if let Some(expr) = level { - self.put_expression(expr, context)?; + self.put_expression(expr, context, true)?; } write!(self.out, ")")?; Ok(()) @@ -209,6 +209,7 @@ impl Writer { &mut self, expr_handle: Handle, context: &ExpressionContext, + is_scoped: bool, ) -> Result<(), Error> { // Add to the set in order to track the stack size. #[cfg(test)] @@ -225,13 +226,13 @@ impl Writer { log::trace!("expression {:?} = {:?}", expr_handle, expression); match *expression { crate::Expression::Access { base, index } => { - self.put_expression(base, context)?; + self.put_expression(base, context, false)?; write!(self.out, "[")?; - self.put_expression(index, context)?; + self.put_expression(index, context, true)?; write!(self.out, "]")?; } crate::Expression::AccessIndex { base, index } => { - self.put_expression(base, context)?; + self.put_expression(base, context, false)?; let base_res = &context.info[base].ty; let mut resolved = base_res.inner_with(&context.module.types); let base_ty_handle = match *resolved { @@ -300,7 +301,7 @@ impl Writer { if i != 0 { write!(self.out, ", ")?; } - self.put_expression(component, context)?; + self.put_expression(component, context, true)?; } write!(self.out, "}}")?; } @@ -335,7 +336,7 @@ impl Writer { } crate::Expression::Load { pointer } => { //write!(self.out, "*")?; - self.put_expression(pointer, context)?; + self.put_expression(pointer, context, is_scoped)?; } crate::Expression::ImageSample { image, @@ -350,18 +351,18 @@ impl Writer { Some(_) => "sample_compare", None => "sample", }; - self.put_expression(image, context)?; + self.put_expression(image, context, false)?; write!(self.out, ".{}(", op)?; - self.put_expression(sampler, context)?; + self.put_expression(sampler, context, true)?; write!(self.out, ", ")?; - self.put_expression(coordinate, context)?; + self.put_expression(coordinate, context, true)?; if let Some(expr) = array_index { write!(self.out, ", ")?; - self.put_expression(expr, context)?; + self.put_expression(expr, context, true)?; } if let Some(dref) = depth_ref { write!(self.out, ", ")?; - self.put_expression(dref, context)?; + self.put_expression(dref, context, true)?; } match level { crate::SampleLevel::Auto => {} @@ -370,19 +371,19 @@ impl Writer { } crate::SampleLevel::Exact(h) => { write!(self.out, ", level(")?; - self.put_expression(h, context)?; + self.put_expression(h, context, true)?; write!(self.out, ")")?; } crate::SampleLevel::Bias(h) => { write!(self.out, ", bias(")?; - self.put_expression(h, context)?; + self.put_expression(h, context, true)?; write!(self.out, ")")?; } crate::SampleLevel::Gradient { x, y } => { write!(self.out, ", gradient(")?; - self.put_expression(x, context)?; + self.put_expression(x, context, true)?; write!(self.out, ", ")?; - self.put_expression(y, context)?; + self.put_expression(y, context, true)?; write!(self.out, ")")?; } } @@ -398,16 +399,16 @@ impl Writer { array_index, index, } => { - self.put_expression(image, context)?; + self.put_expression(image, context, false)?; write!(self.out, ".read(")?; - self.put_expression(coordinate, context)?; + self.put_expression(coordinate, context, true)?; if let Some(expr) = array_index { write!(self.out, ", ")?; - self.put_expression(expr, context)?; + self.put_expression(expr, context, true)?; } if let Some(index) = index { write!(self.out, ", ")?; - self.put_expression(index, context)?; + self.put_expression(index, context, true)?; } write!(self.out, ")")?; } @@ -452,17 +453,17 @@ impl Writer { } crate::ImageQuery::NumLevels => { write!(self.out, "int(")?; - self.put_expression(image, context)?; + self.put_expression(image, context, false)?; write!(self.out, ".get_num_mip_levels())")?; } crate::ImageQuery::NumLayers => { write!(self.out, "int(")?; - self.put_expression(image, context)?; + self.put_expression(image, context, false)?; write!(self.out, ".get_array_size())")?; } crate::ImageQuery::NumSamples => { write!(self.out, "int(")?; - self.put_expression(image, context)?; + self.put_expression(image, context, false)?; write!(self.out, ".get_num_samples())")?; } }, @@ -472,7 +473,7 @@ impl Writer { crate::UnaryOperator::Not => "!", }; write!(self.out, "{}", op_str)?; - self.put_expression(expr, context)?; + self.put_expression(expr, context, false)?; } crate::Expression::Binary { op, left, right } => { let op_str = match op { @@ -501,16 +502,20 @@ impl Writer { .ok_or(Error::UnsupportedBinaryOp(op))?; if op == crate::BinaryOperator::Modulo && kind == crate::ScalarKind::Float { write!(self.out, "fmod(")?; - self.put_expression(left, context)?; + self.put_expression(left, context, true)?; write!(self.out, ", ")?; - self.put_expression(right, context)?; + self.put_expression(right, context, true)?; write!(self.out, ")")?; } else { - write!(self.out, "(")?; - self.put_expression(left, context)?; + if !is_scoped { + write!(self.out, "(")?; + } + self.put_expression(left, context, false)?; write!(self.out, " {} ", op_str)?; - self.put_expression(right, context)?; - write!(self.out, ")")?; + self.put_expression(right, context, false)?; + if !is_scoped { + write!(self.out, ")")?; + } } } crate::Expression::Select { @@ -518,13 +523,17 @@ impl Writer { accept, reject, } => { - write!(self.out, "(")?; - self.put_expression(condition, context)?; + if !is_scoped { + write!(self.out, "(")?; + } + self.put_expression(condition, context, false)?; write!(self.out, " ? ")?; - self.put_expression(accept, context)?; + self.put_expression(accept, context, false)?; write!(self.out, " : ")?; - self.put_expression(reject, context)?; - write!(self.out, ")")?; + self.put_expression(reject, context, false)?; + if !is_scoped { + write!(self.out, ")")?; + } } crate::Expression::Derivative { axis, expr } => { let op = match axis { @@ -628,7 +637,7 @@ impl Writer { }; let op = if convert { "static_cast" } else { "as_type" }; write!(self.out, "{}<{}{}>(", op, scalar, size)?; - self.put_expression(expr, context)?; + self.put_expression(expr, context, true)?; write!(self.out, ")")?; } // has to be a named expression @@ -701,7 +710,7 @@ impl Writer { if min_ref_count <= context.expression.info[handle].ref_count { write!(self.out, "{}", level)?; self.start_baking_expression(handle, &context.expression)?; - self.put_expression(handle, &context.expression)?; + self.put_expression(handle, &context.expression, true)?; writeln!(self.out, ";")?; self.named_expressions.insert(handle.index()); } @@ -720,7 +729,7 @@ impl Writer { ref reject, } => { write!(self.out, "{}if (", level)?; - self.put_expression(condition, &context.expression)?; + self.put_expression(condition, &context.expression, true)?; writeln!(self.out, ") {{")?; self.put_block(level.next(), accept, context)?; if !reject.is_empty() { @@ -735,7 +744,7 @@ impl Writer { ref default, } => { write!(self.out, "{}switch(", level)?; - self.put_expression(selector, &context.expression)?; + self.put_expression(selector, &context.expression, true)?; writeln!(self.out, ") {{")?; let lcase = level.next(); for case in cases.iter() { @@ -789,7 +798,7 @@ impl Writer { } => { let tmp = "_tmp"; write!(self.out, "{}const auto {} = ", level, tmp)?; - self.put_expression(expr_handle, &context.expression)?; + self.put_expression(expr_handle, &context.expression, true)?; writeln!(self.out, ";")?; write!(self.out, "{}return {} {{", level, struct_name)?; for index in 0..members.len() as u32 { @@ -801,14 +810,14 @@ impl Writer { } _ => { write!(self.out, "{}return {} {{ ", level, struct_name)?; - self.put_expression(expr_handle, &context.expression)?; + self.put_expression(expr_handle, &context.expression, true)?; } } write!(self.out, " }}")?; } None => { write!(self.out, "{}return ", level)?; - self.put_expression(expr_handle, &context.expression)?; + self.put_expression(expr_handle, &context.expression, true)?; } } writeln!(self.out, ";")?; @@ -821,9 +830,9 @@ impl Writer { } crate::Statement::Store { pointer, value } => { write!(self.out, "{}", level)?; - self.put_expression(pointer, &context.expression)?; + self.put_expression(pointer, &context.expression, true)?; write!(self.out, " = ")?; - self.put_expression(value, &context.expression)?; + self.put_expression(value, &context.expression, true)?; writeln!(self.out, ";")?; } crate::Statement::ImageStore { @@ -833,14 +842,14 @@ impl Writer { value, } => { write!(self.out, "{}", level)?; - self.put_expression(image, &context.expression)?; + self.put_expression(image, &context.expression, false)?; write!(self.out, ".write(")?; - self.put_expression(value, &context.expression)?; + self.put_expression(value, &context.expression, true)?; write!(self.out, ", ")?; - self.put_expression(coordinate, &context.expression)?; + self.put_expression(coordinate, &context.expression, true)?; if let Some(expr) = array_index { write!(self.out, ", ")?; - self.put_expression(expr, &context.expression)?; + self.put_expression(expr, &context.expression, true)?; } writeln!(self.out, ");")?; } @@ -861,7 +870,7 @@ impl Writer { if i != 0 { write!(self.out, ", ")?; } - self.put_expression(handle, &context.expression)?; + self.put_expression(handle, &context.expression, true)?; } // follow-up with any global resources used let mut separate = !arguments.is_empty(); @@ -1513,8 +1522,8 @@ fn test_stack_size() { } let stack_size = max_addr - min_addr; // check the size (in debug only) - // last observed macOS value: 20672 - if stack_size > 21000 { + // last observed macOS value: 21920 + if stack_size > 22000 { panic!("`put_expression` stack size {} is too large!", stack_size); } } diff --git a/tests/out/boids.msl.snap b/tests/out/boids.msl.snap index 816ad79a49..a7d8102413 100644 --- a/tests/out/boids.msl.snap +++ b/tests/out/boids.msl.snap @@ -54,7 +54,7 @@ kernel void main1( type1 pos1; type1 vel1; type i = const_0u; - if ((global_invocation_id.x >= NUM_PARTICLES)) { + if (global_invocation_id.x >= NUM_PARTICLES) { return; } vPos = particlesSrc.particles[global_invocation_id.x].pos; @@ -65,48 +65,48 @@ kernel void main1( bool loop_init = true; while(true) { if (!loop_init) { - i = (i + const_1u); + i = i + const_1u; } loop_init = false; - if ((i >= NUM_PARTICLES)) { + if (i >= NUM_PARTICLES) { break; } - if ((i == global_invocation_id.x)) { + if (i == global_invocation_id.x) { continue; } pos1 = particlesSrc.particles[i].pos; vel1 = particlesSrc.particles[i].vel; - if ((metal::distance(pos1, vPos) < params.rule1Distance)) { - cMass = (cMass + pos1); - cMassCount = (cMassCount + const_1i); + if (metal::distance(pos1, vPos) < params.rule1Distance) { + cMass = cMass + pos1; + cMassCount = cMassCount + const_1i; } - if ((metal::distance(pos1, vPos) < params.rule2Distance)) { - colVel = (colVel - (pos1 - vPos)); + if (metal::distance(pos1, vPos) < params.rule2Distance) { + colVel = colVel - (pos1 - vPos); } - if ((metal::distance(pos1, vPos) < params.rule3Distance)) { - cVel = (cVel + vel1); - cVelCount = (cVelCount + const_1i); + if (metal::distance(pos1, vPos) < params.rule3Distance) { + cVel = cVel + vel1; + cVelCount = cVelCount + const_1i; } } - if ((cMassCount > const_0i)) { - cMass = ((cMass * (const_1f / static_cast(cMassCount))) - vPos); + if (cMassCount > const_0i) { + cMass = (cMass * (const_1f / static_cast(cMassCount))) - vPos; } - if ((cVelCount > const_0i)) { - cVel = (cVel * (const_1f / static_cast(cVelCount))); + if (cVelCount > const_0i) { + cVel = cVel * (const_1f / static_cast(cVelCount)); } - vVel = (((vVel + (cMass * params.rule1Scale)) + (colVel * params.rule2Scale)) + (cVel * params.rule3Scale)); - vVel = (metal::normalize(vVel) * metal::clamp(metal::length(vVel), const_0f, const_0_10f)); - vPos = (vPos + (vVel * params.deltaT)); - if ((vPos.x < const_n1f)) { + vVel = ((vVel + (cMass * params.rule1Scale)) + (colVel * params.rule2Scale)) + (cVel * params.rule3Scale); + vVel = metal::normalize(vVel) * metal::clamp(metal::length(vVel), const_0f, const_0_10f); + vPos = vPos + (vVel * params.deltaT); + if (vPos.x < const_n1f) { vPos.x = const_1f; } - if ((vPos.x > const_1f)) { + if (vPos.x > const_1f) { vPos.x = const_n1f; } - if ((vPos.y < const_n1f)) { + if (vPos.y < const_n1f) { vPos.y = const_1f; } - if ((vPos.y > const_1f)) { + if (vPos.y > const_1f) { vPos.y = const_n1f; } particlesDst.particles[global_invocation_id.x].pos = vPos; diff --git a/tests/out/collatz.msl.snap b/tests/out/collatz.msl.snap index 6c9e2f0a65..b7d56f09c2 100644 --- a/tests/out/collatz.msl.snap +++ b/tests/out/collatz.msl.snap @@ -22,15 +22,15 @@ type collatz_iterations( type i = const_0u; n = n_base; while(true) { - if ((n <= const_1u)) { + if (n <= const_1u) { break; } - if (((n % const_2u) == const_0u)) { - n = (n / const_2u); + if ((n % const_2u) == const_0u) { + n = n / const_2u; } else { - n = ((const_3u * n) + const_1u); + n = (const_3u * n) + const_1u; } - i = (i + const_1u); + i = i + const_1u; } return i; } diff --git a/tests/out/quad.msl.snap b/tests/out/quad.msl.snap index ebea304a27..0a01f1484a 100644 --- a/tests/out/quad.msl.snap +++ b/tests/out/quad.msl.snap @@ -32,7 +32,7 @@ vertex main1Output main1( const auto uv1 = varyings.uv1; VertexOutput out; out.uv = uv1; - out.position = metal::float4((c_scale * pos), const_0f, const_1f); + out.position = metal::float4(c_scale * pos, const_0f, const_1f); const auto _tmp = out; return main1Output { _tmp.uv, _tmp.position }; } @@ -50,9 +50,9 @@ fragment main2Output main2( ) { const auto uv2 = varyings1.uv2; metal::float4 _expr4 = u_texture.sample(u_sampler, uv2); - if ((_expr4.w == const_0f)) { + if (_expr4.w == const_0f) { metal::discard_fragment(); } - return main2Output { (_expr4.w * _expr4) }; + return main2Output { _expr4.w * _expr4 }; } diff --git a/tests/out/shadow.msl.snap b/tests/out/shadow.msl.snap index 0e2e670670..b2eae0b4c7 100644 --- a/tests/out/shadow.msl.snap +++ b/tests/out/shadow.msl.snap @@ -41,11 +41,11 @@ type7 fetch_shadow( type4 t_shadow, type5 sampler_shadow ) { - if ((homogeneous_coords.w <= const_0f)) { + if (homogeneous_coords.w <= const_0f) { return const_1f; } - float _expr15 = (const_1f / homogeneous_coords.w); - float _expr28 = t_shadow.sample_compare(sampler_shadow, (((metal::float2(homogeneous_coords.x, homogeneous_coords.y) * metal::float2(const_0_50f, const_n0_50f)) * _expr15) + metal::float2(const_0_50f, const_0_50f)), static_cast(light_id), (homogeneous_coords.z * _expr15)); + float _expr15 = const_1f / homogeneous_coords.w; + float _expr28 = t_shadow.sample_compare(sampler_shadow, ((metal::float2(homogeneous_coords.x, homogeneous_coords.y) * metal::float2(const_0_50f, const_n0_50f)) * _expr15) + metal::float2(const_0_50f, const_0_50f), static_cast(light_id), homogeneous_coords.z * _expr15); return _expr28; } @@ -70,15 +70,15 @@ fragment fs_mainOutput fs_main( bool loop_init = true; while(true) { if (!loop_init) { - i = (i + const_1u); + i = i + const_1u; } loop_init = false; - if ((i >= metal::min(u_globals.num_lights.x, c_max_lights))) { + if (i >= metal::min(u_globals.num_lights.x, c_max_lights)) { break; } Light _expr21 = s_lights.data[i]; - type7 _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))); + type7 _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 83ce8194a4..eb27b7cfeb 100644 --- a/tests/out/skybox.msl.snap +++ b/tests/out/skybox.msl.snap @@ -39,11 +39,11 @@ vertex vs_mainOutput vs_main( type4 tmp1_; type4 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 _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)); + 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 _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; const auto _tmp = out; return vs_mainOutput { _tmp.position, _tmp.uv };