nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-11-26 16:53:48 +00:00
Code Issues Packages Projects Releases Wiki Activity

[msl-out] Improve formatting of whitespace

Browse Source
This commit is contained in:
Joshua Groves 2021-01-25 20:30:23 -03:30 committed by Dzmitry Malyshau
parent a5184fba49
commit 58f38ba61d
6 changed files with 253 additions and 193 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

51
src/back/msl/writer.rs
View File

@ -10,6 +10,7 @@ use std::{
};
const NAMESPACE: &str = "metal";
const INDENT: &str = " ";
struct Level(usize);
impl Level {
@ -19,7 +20,7 @@ impl Level {
}
impl Display for Level {
fn fmt(&self, formatter: &mut Formatter<'_>) -> Result<(), FmtError> {
(0..self.0).try_for_each(|_| formatter.write_str("\t"))
(0..self.0).try_for_each(|_| formatter.write_str(INDENT))
}
}
@ -649,7 +650,7 @@ impl<W: Write> Writer<W> {
writeln!(self.out, "{}discard_fragment();", level)?;
}
crate::Statement::Store { pointer, value } => {
//write!(self.out, "\t*")?;
//write!(self.out, "{}*", INDENT)?;
write!(self.out, "{}", level)?;
self.put_expression(pointer, context)?;
write!(self.out, " = ")?;
@ -671,11 +672,9 @@ impl<W: Write> Writer<W> {
writeln!(self.out, "#include <metal_stdlib>")?;
writeln!(self.out, "#include <simd/simd.h>")?;
writeln!(self.out)?;
self.write_type_defs(module)?;
writeln!(self.out)?;
self.write_functions(module, options)
}
@ -745,7 +744,7 @@ impl<W: Write> Writer<W> {
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)];
writeln!(self.out, "\t{} {};", base_name, member_name)?;
writeln!(self.out, "{}{} {};", INDENT, base_name, member_name)?;
}
write!(self.out, "}}")?;
}
@ -809,6 +808,7 @@ impl<W: Write> Writer<W> {
}
}
writeln!(self.out, ";")?;
writeln!(self.out)?;
}
Ok(())
}
@ -843,14 +843,18 @@ impl<W: Write> Writer<W> {
let name = &self.names[&NameKey::FunctionArgument(fun_handle, index as u32)];
let param_type_name = &self.names[&NameKey::Type(arg.ty)];
let separator = separate(index + 1 == fun.arguments.len());
writeln!(self.out, "\t{} {}{}", param_type_name, name, separator)?;
writeln!(
self.out,
"{}{} {}{}",
INDENT, param_type_name, name, separator
)?;
}
writeln!(self.out, ") {{")?;
for (local_handle, local) in fun.local_variables.iter() {
let ty_name = &self.names[&NameKey::Type(local.ty)];
let local_name = &self.names[&NameKey::FunctionLocal(fun_handle, local_handle)];
write!(self.out, "\t{} {}", ty_name, local_name)?;
write!(self.out, "{}{} {}", INDENT, ty_name, local_name)?;
if let Some(value) = local.init {
write!(self.out, " = ")?;
self.put_constant(value, module)?;
@ -865,6 +869,7 @@ impl<W: Write> Writer<W> {
};
self.put_block(Level(1), &fun.body, &context, None)?;
writeln!(self.out, "}}")?;
writeln!(self.out)?;
}
let mut info = TranslationInfo {
@ -944,12 +949,14 @@ impl<W: Write> Writer<W> {
handle,
usage: crate::GlobalUse::empty(),
};
write!(self.out, "\t")?;
write!(self.out, "{}", INDENT)?;
tyvar.try_fmt(&mut self.out)?;
let resolved = options.resolve_binding(stage, var, in_mode)?;
resolved.try_fmt_decorated(&mut self.out, ";\n")?;
resolved.try_fmt_decorated(&mut self.out, ";")?;
writeln!(self.out)?;
}
writeln!(self.out, "}};")?;
writeln!(self.out)?;
writeln!(self.out, "struct {} {{", output_name)?;
for ((handle, var), &usage) in
@ -967,19 +974,21 @@ impl<W: Write> Writer<W> {
handle,
usage: crate::GlobalUse::empty(),
};
write!(self.out, "\t")?;
write!(self.out, "{}", INDENT)?;
tyvar.try_fmt(&mut self.out)?;
let resolved = options.resolve_binding(stage, var, out_mode)?;
resolved.try_fmt_decorated(&mut self.out, ";\n")?;
resolved.try_fmt_decorated(&mut self.out, ";")?;
writeln!(self.out)?;
}
writeln!(self.out, "}};")?;
writeln!(self.out)?;
writeln!(self.out, "{} {} {}(", em_str, output_name, fun_name)?;
let separator = separate(last_used_global.is_none());
writeln!(
self.out,
"\t{} {} [[stage_in]]{}",
location_input_name, LOCATION_INPUT_STRUCT_NAME, separator
"{}{} {} [[stage_in]]{}",
INDENT, location_input_name, LOCATION_INPUT_STRUCT_NAME, separator
)?;
Some(OUTPUT_STRUCT_NAME)
@ -1021,7 +1030,7 @@ impl<W: Write> Writer<W> {
usage,
};
let separator = separate(last_used_global == Some(handle));
write!(self.out, "\t")?;
write!(self.out, "{}", INDENT)?;
tyvar.try_fmt(&mut self.out)?;
resolved.try_fmt_decorated(&mut self.out, separator)?;
if let Some(value) = var.init {
@ -1034,14 +1043,18 @@ impl<W: Write> Writer<W> {
match stage {
crate::ShaderStage::Vertex | crate::ShaderStage::Fragment => {
writeln!(self.out, "\t{} {};", output_name, OUTPUT_STRUCT_NAME)?;
writeln!(
self.out,
"{}{} {};",
INDENT, output_name, OUTPUT_STRUCT_NAME
)?;
}
crate::ShaderStage::Compute => {}
}
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)];
write!(self.out, "\t{} {}", ty_name, name)?;
write!(self.out, "{}{} {}", INDENT, ty_name, name)?;
if let Some(value) = local.init {
write!(self.out, " = ")?;
self.put_constant(value, module)?;
@ -1056,6 +1069,10 @@ impl<W: Write> Writer<W> {
};
self.put_block(Level(1), &fun.body, &context, return_value)?;
writeln!(self.out, "}}")?;
let is_last = ep_index == module.entry_points.len() - 1;
if !is_last {
writeln!(self.out)?;
}
}
Ok(info)

211
tests/snapshots/snapshots__boids.msl.snap
View File

@ -6,122 +6,137 @@ expression: msl
#include <simd/simd.h>
typedef metal::float2 type;
typedef metal::float4 type1;
typedef float type2;
struct Particle {
type pos;
type vel;
type pos;
type vel;
};
struct SimParams {
type2 deltaT;
type2 rule1Distance;
type2 rule2Distance;
type2 rule3Distance;
type2 rule1Scale;
type2 rule2Scale;
type2 rule3Scale;
type2 deltaT;
type2 rule1Distance;
type2 rule2Distance;
type2 rule3Distance;
type2 rule1Scale;
type2 rule2Scale;
type2 rule3Scale;
};
typedef Particle type3[5];
struct Particles {
type3 particles;
type3 particles;
};
typedef metal::uint3 type4;
typedef uint type5;
typedef int type6;
struct main1Input {
type a_particlePos [[attribute(0)]];
type a_particleVel [[attribute(1)]];
type a_pos [[attribute(2)]];
type a_particlePos [[attribute(0)]];
type a_particleVel [[attribute(1)]];
type a_pos [[attribute(2)]];
};
struct main1Output {
type1 gl_Position [[position]];
type1 gl_Position [[position]];
};
vertex main1Output main1(
main1Input input [[stage_in]]
main1Input input [[stage_in]]
) {
main1Output output;
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 {
};
struct main2Output {
type1 fragColor [[color(0)]];
};
fragment main2Output main2(
main2Input input [[stage_in]]
) {
main2Output output;
output.fragColor = metal::float4(1.0, 1.0, 1.0, 1.0);
return output;
}
kernel void main3(
constant SimParams& params [[buffer(0)]],
constant Particles& particlesA [[buffer(1)]],
device Particles& particlesB [[buffer(2)]],
type4 gl_GlobalInvocationID [[thread_position_in_grid]]
) {
type vPos;
type vVel;
type cMass;
type cVel;
type colVel;
type6 cMassCount = 0;
type6 cVelCount = 0;
type pos1;
type vel1;
type5 i = 0;
if (gl_GlobalInvocationID.x >= 5) {
}
vPos = particlesA.particles[gl_GlobalInvocationID.x].pos;
vVel = particlesA.particles[gl_GlobalInvocationID.x].vel;
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 >= 5) {
break;
}
if (i == gl_GlobalInvocationID.x) {
continue;
}
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;
}
if (metal::distance(pos1, vPos) < params.rule2Distance) {
colVel = colVel - pos1 - vPos;
}
if (metal::distance(pos1, vPos) < params.rule3Distance) {
cVel = cVel + vel1;
cVelCount = cVelCount + 1;
}
}
if (cMassCount == 0) {
cMass = cMass / metal::float2(cMassCount, cMassCount) + vPos;
}
if (cVelCount == 0) {
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);
vPos = vPos + vVel * params.deltaT;
if (vPos.x < -1.0) {
vPos.x = 1.0;
}
if (vPos.x == 1.0) {
vPos.x = -1.0;
}
if (vPos.y < -1.0) {
vPos.y = 1.0;
}
if (vPos.y == 1.0) {
vPos.y = -1.0;
}
particlesB.particles[gl_GlobalInvocationID.x].pos = vPos;
particlesB.particles[gl_GlobalInvocationID.x].vel = vVel;
main1Output output;
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 {
};
struct main2Output {
type1 fragColor [[color(0)]];
};
fragment main2Output main2(
main2Input input [[stage_in]]
) {
main2Output output;
output.fragColor = metal::float4(1.0, 1.0, 1.0, 1.0);
return output;
}
kernel void main3(
constant SimParams& params [[buffer(0)]],
constant Particles& particlesA [[buffer(1)]],
device Particles& particlesB [[buffer(2)]],
type4 gl_GlobalInvocationID [[thread_position_in_grid]]
) {
type vPos;
type vVel;
type cMass;
type cVel;
type colVel;
type6 cMassCount = 0;
type6 cVelCount = 0;
type pos1;
type vel1;
type5 i = 0;
if (gl_GlobalInvocationID.x >= 5) {
}
vPos = particlesA.particles[gl_GlobalInvocationID.x].pos;
vVel = particlesA.particles[gl_GlobalInvocationID.x].vel;
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 >= 5) {
break;
}
if (i == gl_GlobalInvocationID.x) {
continue;
}
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;
}
if (metal::distance(pos1, vPos) < params.rule2Distance) {
colVel = colVel - pos1 - vPos;
}
if (metal::distance(pos1, vPos) < params.rule3Distance) {
cVel = cVel + vel1;
cVelCount = cVelCount + 1;
}
}
if (cMassCount == 0) {
cMass = cMass / metal::float2(cMassCount, cMassCount) + vPos;
}
if (cVelCount == 0) {
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);
vPos = vPos + vVel * params.deltaT;
if (vPos.x < -1.0) {
vPos.x = 1.0;
}
if (vPos.x == 1.0) {
vPos.x = -1.0;
}
if (vPos.y < -1.0) {
vPos.y = 1.0;
}
if (vPos.y == 1.0) {
vPos.y = -1.0;
}
particlesB.particles[gl_GlobalInvocationID.x].pos = vPos;
particlesB.particles[gl_GlobalInvocationID.x].vel = vVel;
}

46
tests/snapshots/snapshots__collatz.msl.snap
View File

@ -6,33 +6,37 @@ expression: msl
#include <simd/simd.h>
typedef metal::uint3 type;
typedef uint type1;
typedef type1 type2[1];
struct PrimeIndices {
type2 data;
type2 data;
};
type1 collatz_iterations(
type1 n
type1 n
) {
type1 i = 0;
while(true) {
if (n <= 1) {
break;
}
if (n % 2 == 0) {
n = n / 2;
} else {
n = 3 * n + 1;
}
i = i + 1;
}
return i;
}
kernel void main1(
type global_id [[thread_position_in_grid]],
device PrimeIndices& v_indices [[buffer(0)]]
) {
v_indices.data[global_id.x] = collatz_iterations(v_indices.data[global_id.x]);
type1 i = 0;
while(true) {
if (n <= 1) {
break;
}
if (n % 2 == 0) {
n = n / 2;
} else {
n = 3 * n + 1;
}
i = i + 1;
}
return i;
}
kernel void main1(
type global_id [[thread_position_in_grid]],
device PrimeIndices& v_indices [[buffer(0)]]
) {
v_indices.data[global_id.x] = collatz_iterations(v_indices.data[global_id.x]);
}

1
tests/snapshots/snapshots__empty.msl.snap
View File

@ -5,7 +5,6 @@ expression: msl
#include <metal_stdlib>
#include <simd/simd.h>
kernel void main1(
) {
}

58
tests/snapshots/snapshots__quad.msl.snap
View File

@ -6,41 +6,51 @@ expression: msl
#include <simd/simd.h>
typedef float type;
typedef metal::float2 type1;
typedef metal::float4 type2;
typedef metal::texture2d<float, metal::access::sample> type3;
typedef metal::sampler type4;
typedef metal::int2 type5;
struct main1Input {
type1 a_pos [[attribute(0)]];
type1 a_uv [[attribute(1)]];
type1 a_pos [[attribute(0)]];
type1 a_uv [[attribute(1)]];
};
struct main1Output {
type1 v_uv [[user(loc0)]];
type2 o_position [[position]];
type1 v_uv [[user(loc0)]];
type2 o_position [[position]];
};
vertex main1Output main1(
main1Input input [[stage_in]]
main1Input input [[stage_in]]
) {
main1Output output;
output.v_uv = input.a_uv;
output.o_position = metal::float4(1.2 * input.a_pos, 0.0, 1.0);
return output;
}
struct main2Input {
type1 v_uv1 [[user(loc0)]];
};
struct main2Output {
type2 o_color [[color(0)]];
};
fragment main2Output main2(
main2Input input [[stage_in]],
type3 u_texture [[texture(0)]],
type4 u_sampler [[sampler(0)]]
) {
main2Output output;
output.o_color = u_texture.sample(u_sampler, input.v_uv1);
return output;
main1Output output;
output.v_uv = input.a_uv;
output.o_position = metal::float4(1.2 * input.a_pos, 0.0, 1.0);
return output;
}
struct main2Input {
type1 v_uv1 [[user(loc0)]];
};
struct main2Output {
type2 o_color [[color(0)]];
};
fragment main2Output main2(
main2Input input [[stage_in]],
type3 u_texture [[texture(0)]],
type4 u_sampler [[sampler(0)]]
) {
main2Output output;
output.o_color = u_texture.sample(u_sampler, input.v_uv1);
return output;
}

79
tests/snapshots/snapshots__skybox.msl.snap
View File

@ -6,55 +6,70 @@ expression: msl
#include <simd/simd.h>
typedef metal::float4 type;
typedef metal::float3 type1;
typedef uint type2;
typedef metal::float4x4 type3;
struct Data {
type3 proj_inv;
type3 view;
type3 proj_inv;
type3 view;
};
typedef int type4;
typedef metal::float3x3 type5;
typedef float type6;
typedef metal::texturecube<float, metal::access::sample> type7;
typedef metal::sampler type8;
typedef metal::int3 type9;
struct vs_mainInput {
};
struct vs_mainOutput {
type out_position [[position]];
type1 out_uv [[user(loc0)]];
type out_position [[position]];
type1 out_uv [[user(loc0)]];
};
vertex vs_mainOutput vs_main(
vs_mainInput input [[stage_in]],
type2 in_vertex_index [[vertex_id]],
constant Data& r_data [[buffer(0)]]
vs_mainInput input [[stage_in]],
type2 in_vertex_index [[vertex_id]],
constant Data& r_data [[buffer(0)]]
) {
vs_mainOutput output;
type4 tmp1_;
type4 tmp2_;
type unprojected;
tmp1_ = static_cast<int>(in_vertex_index) / 2;
tmp2_ = static_cast<int>(in_vertex_index) & 1;
unprojected = r_data.proj_inv * metal::float4(static_cast<float>(tmp1_) * 4.0 - 1.0, static_cast<float>(tmp2_) * 4.0 - 1.0, 0.0, 1.0);
output.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(unprojected.x, unprojected.y, unprojected.z);
output.out_position = metal::float4(static_cast<float>(tmp1_) * 4.0 - 1.0, static_cast<float>(tmp2_) * 4.0 - 1.0, 0.0, 1.0);
return output;
}
struct fs_mainInput {
type1 in_uv [[user(loc0)]];
};
struct fs_mainOutput {
type out_color [[color(0)]];
};
fragment fs_mainOutput fs_main(
fs_mainInput input [[stage_in]],
type7 r_texture [[texture(0)]],
type8 r_sampler [[sampler(1)]]
) {
fs_mainOutput output;
output.out_color = r_texture.sample(r_sampler, input.in_uv);
return output;
vs_mainOutput output;
type4 tmp1_;
type4 tmp2_;
type unprojected;
tmp1_ = static_cast<int>(in_vertex_index) / 2;
tmp2_ = static_cast<int>(in_vertex_index) & 1;
unprojected = r_data.proj_inv * metal::float4(static_cast<float>(tmp1_) * 4.0 - 1.0, static_cast<float>(tmp2_) * 4.0 - 1.0, 0.0, 1.0);
output.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(unprojected.x, unprojected.y, unprojected.z);
output.out_position = metal::float4(static_cast<float>(tmp1_) * 4.0 - 1.0, static_cast<float>(tmp2_) * 4.0 - 1.0, 0.0, 1.0);
return output;
}
struct fs_mainInput {
type1 in_uv [[user(loc0)]];
};
struct fs_mainOutput {
type out_color [[color(0)]];
};
fragment fs_mainOutput fs_main(
fs_mainInput input [[stage_in]],
type7 r_texture [[texture(0)]],
type8 r_sampler [[sampler(1)]]
) {
fs_mainOutput output;
output.out_color = r_texture.sample(r_sampler, input.in_uv);
return output;
}