[msl] don't put expressions in double scopes

This commit is contained in:
Dzmitry Malyshau
2021-03-26 00:14:01 -04:00
committed by Dzmitry Malyshau
parent 2b163dc2a1
commit bf193bb535
6 changed files with 104 additions and 95 deletions

View File

@@ -183,7 +183,7 @@ impl<W: Write> Writer<W> {
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<W: Write> Writer<W> {
level: Option<Handle<crate::Expression>>,
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<W: Write> Writer<W> {
&mut self,
expr_handle: Handle<crate::Expression>,
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
}
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
}
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
}
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
.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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
};
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
} => {
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<W: Write> Writer<W> {
}
_ => {
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<W: Write> Writer<W> {
}
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
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);
}
}

View File

@@ -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<float>(cMassCount))) - vPos);
if (cMassCount > const_0i) {
cMass = (cMass * (const_1f / static_cast<float>(cMassCount))) - vPos;
}
if ((cVelCount > const_0i)) {
cVel = (cVel * (const_1f / static_cast<float>(cVelCount)));
if (cVelCount > const_0i) {
cVel = cVel * (const_1f / static_cast<float>(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;

View File

@@ -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;
}

View File

@@ -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 };
}

View File

@@ -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<int>(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<int>(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) };
}

View File

@@ -39,11 +39,11 @@ vertex vs_mainOutput vs_main(
type4 tmp1_;
type4 tmp2_;
VertexOutput out;
tmp1_ = (static_cast<int>(vertex_index) / const_2i);
tmp2_ = (static_cast<int>(vertex_index) & const_1i);
type _expr24 = metal::float4(((static_cast<float>(tmp1_) * const_4f) - const_1f), ((static_cast<float>(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<int>(vertex_index) / const_2i;
tmp2_ = static_cast<int>(vertex_index) & const_1i;
type _expr24 = metal::float4((static_cast<float>(tmp1_) * const_4f) - const_1f, (static_cast<float>(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 };