github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

[msl] fix metal namespacing, matrix construction and indexing

Browse Source
This commit is contained in:
Dzmitry Malyshau
2020-12-12 13:01:25 -05:00
committed by Dzmitry Malyshau
parent 933ca3863b
commit 018dfd6d29
4 changed files with 58 additions and 30 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
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<W: Write> Writer<W> {
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, ";")?;

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

@@ -5,8 +5,8 @@ expression: msl
#include <metal_stdlib>
#include <simd/simd.h>
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<uint>(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);

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

@@ -6,11 +6,11 @@ expression: msl
#include <simd/simd.h>
typedef float type;
typedef float2 type1;
typedef float4 type2;
typedef texture2d<float, access::sample> type3;
typedef sampler type4;
typedef int2 type5;
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)]];
@@ -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 {

4
tests/snapshots/snapshots__simple.msl.snap
View File

@@ -5,7 +5,7 @@ expression: msl
#include <metal_stdlib>
#include <simd/simd.h>
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;
}