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] write scalar constants before the types

Browse Source
This commit is contained in:
Dzmitry Malyshau
2021-03-26 00:21:17 -04:00
committed by Dzmitry Malyshau
parent bf193bb535
commit a090afa3f5
6 changed files with 70 additions and 52 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -501,7 +501,7 @@ impl<W: Write> Writer<W> {
.scalar_kind()
.ok_or(Error::UnsupportedBinaryOp(op))?;
if op == crate::BinaryOperator::Modulo && kind == crate::ScalarKind::Float {
write!(self.out, "fmod(")?;
write!(self.out, "{}::fmod(", NAMESPACE)?;
self.put_expression(left, context, true)?;
write!(self.out, ", ")?;
self.put_expression(right, context, true)?;
@@ -907,8 +907,9 @@ impl<W: Write> Writer<W> {
writeln!(self.out, "#include <simd/simd.h>")?;
writeln!(self.out)?;
self.write_scalar_constants(module)?;
self.write_type_defs(module)?;
self.write_constants(module)?;
self.write_composite_constants(module)?;
self.write_functions(module, info, options)
}
@@ -1076,43 +1077,60 @@ impl<W: Write> Writer<W> {
Ok(())
}
fn write_constants(&mut self, module: &crate::Module) -> Result<(), Error> {
fn write_scalar_constants(&mut self, module: &crate::Module) -> Result<(), Error> {
for (handle, constant) in module.constants.iter() {
write!(self.out, "constexpr constant ")?;
let name = &self.names[&NameKey::Constant(handle)];
match constant.inner {
crate::ConstantInner::Scalar {
width: _,
ref value,
} => match *value {
crate::ScalarValue::Sint(value) => {
write!(self.out, "int {} = {}", name, value)?;
}
crate::ScalarValue::Uint(value) => {
write!(self.out, "unsigned {} = {}u", name, value)?;
}
crate::ScalarValue::Float(value) => {
write!(self.out, "float {} = {}", name, value)?;
if value.fract() == 0.0 {
write!(self.out, ".0")?;
} => {
let name = &self.names[&NameKey::Constant(handle)];
write!(self.out, "constexpr constant ")?;
match *value {
crate::ScalarValue::Sint(value) => {
write!(self.out, "int {} = {}", name, value)?;
}
crate::ScalarValue::Uint(value) => {
write!(self.out, "unsigned {} = {}u", name, value)?;
}
crate::ScalarValue::Float(value) => {
write!(self.out, "float {} = {}", name, value)?;
if value.fract() == 0.0 {
write!(self.out, ".0")?;
}
}
crate::ScalarValue::Bool(value) => {
write!(self.out, "bool {} = {}", name, value)?;
}
}
crate::ScalarValue::Bool(value) => {
write!(self.out, "bool {} = {}", name, value)?;
}
},
writeln!(self.out, ";")?;
}
crate::ConstantInner::Composite { .. } => {}
}
}
Ok(())
}
fn write_composite_constants(&mut self, module: &crate::Module) -> Result<(), Error> {
for (handle, constant) in module.constants.iter() {
match constant.inner {
crate::ConstantInner::Scalar { .. } => {}
crate::ConstantInner::Composite { ty, ref components } => {
let name = &self.names[&NameKey::Constant(handle)];
let ty_name = &self.names[&NameKey::Type(ty)];
write!(self.out, "{} {} = {}(", ty_name, name, ty_name)?;
write!(
self.out,
"constexpr constant {} {} = {}(",
ty_name, name, ty_name
)?;
for (i, &sub_handle) in components.iter().enumerate() {
let separator = if i != 0 { ", " } else { "" };
let sub_name = &self.names[&NameKey::Constant(sub_handle)];
write!(self.out, "{}{}", separator, sub_name)?;
}
write!(self.out, ")")?;
writeln!(self.out, ");")?;
}
}
writeln!(self.out, ";")?;
}
Ok(())
}

18
tests/out/boids.msl.snap
View File

@@ -5,6 +5,15 @@ expression: msl
#include <metal_stdlib>
#include <simd/simd.h>
constexpr constant unsigned NUM_PARTICLES = 1500u;
constexpr constant float const_0f = 0.0;
constexpr constant int const_0i = 0;
constexpr constant unsigned const_0u = 0u;
constexpr constant int const_1i = 1;
constexpr constant unsigned const_1u = 1u;
constexpr constant float const_1f = 1.0;
constexpr constant float const_0_10f = 0.1;
constexpr constant float const_n1f = -1.0;
typedef uint type;
typedef metal::float2 type1;
struct Particle {
@@ -27,15 +36,6 @@ struct Particles {
};
typedef metal::uint3 type4;
typedef int type5;
constexpr constant unsigned NUM_PARTICLES = 1500u;
constexpr constant float const_0f = 0.0;
constexpr constant int const_0i = 0;
constexpr constant unsigned const_0u = 0u;
constexpr constant int const_1i = 1;
constexpr constant unsigned const_1u = 1u;
constexpr constant float const_1f = 1.0;
constexpr constant float const_0_10f = 0.1;
constexpr constant float const_n1f = -1.0;
struct main1Input {
};
kernel void main1(

8
tests/out/collatz.msl.snap
View File

@@ -5,16 +5,16 @@ expression: msl
#include <metal_stdlib>
#include <simd/simd.h>
constexpr constant unsigned const_0u = 0u;
constexpr constant unsigned const_1u = 1u;
constexpr constant unsigned const_2u = 2u;
constexpr constant unsigned const_3u = 3u;
typedef uint type;
typedef type type1[1];
struct PrimeIndices {
type1 data;
};
typedef metal::uint3 type2;
constexpr constant unsigned const_0u = 0u;
constexpr constant unsigned const_1u = 1u;
constexpr constant unsigned const_2u = 2u;
constexpr constant unsigned const_3u = 3u;
type collatz_iterations(
type n_base
) {

6
tests/out/quad.msl.snap
View File

@@ -5,6 +5,9 @@ expression: msl
#include <metal_stdlib>
#include <simd/simd.h>
constexpr constant float c_scale = 1.2;
constexpr constant float const_0f = 0.0;
constexpr constant float const_1f = 1.0;
typedef float type;
typedef metal::float2 type1;
typedef metal::float4 type2;
@@ -14,9 +17,6 @@ struct VertexOutput {
};
typedef metal::texture2d<float, metal::access::sample> type3;
typedef metal::sampler type4;
constexpr constant float c_scale = 1.2;
constexpr constant float const_0f = 0.0;
constexpr constant float const_1f = 1.0;
struct main1Input {
type1 pos [[attribute(0)]];
type1 uv1 [[attribute(1)]];

16
tests/out/shadow.msl.snap
View File

@@ -5,6 +5,14 @@ expression: msl
#include <metal_stdlib>
#include <simd/simd.h>
constexpr constant float const_0f = 0.0;
constexpr constant float const_1f = 1.0;
constexpr constant float const_0_50f = 0.5;
constexpr constant float const_n0_50f = -0.5;
constexpr constant float const_0_05f = 0.05;
constexpr constant unsigned c_max_lights = 10u;
constexpr constant unsigned const_0u = 0u;
constexpr constant unsigned const_1u = 1u;
typedef metal::uint4 type;
struct Globals {
type num_lights;
@@ -26,15 +34,7 @@ typedef uint type6;
typedef float type7;
typedef metal::float2 type8;
typedef metal::float3 type9;
constexpr constant float const_0f = 0.0;
constexpr constant float const_1f = 1.0;
constexpr constant float const_0_50f = 0.5;
constexpr constant float const_n0_50f = -0.5;
constexpr constant float const_0_05f = 0.05;
constexpr constant type9 c_ambient = type9(const_0_05f, const_0_05f, const_0_05f);
constexpr constant unsigned c_max_lights = 10u;
constexpr constant unsigned const_0u = 0u;
constexpr constant unsigned const_1u = 1u;
type7 fetch_shadow(
type6 light_id,
type2 homogeneous_coords,

10
tests/out/skybox.msl.snap
View File

@@ -5,6 +5,11 @@ expression: msl
#include <metal_stdlib>
#include <simd/simd.h>
constexpr constant int const_2i = 2;
constexpr constant int const_1i = 1;
constexpr constant float const_4f = 4.0;
constexpr constant float const_1f = 1.0;
constexpr constant float const_0f = 0.0;
typedef metal::float4 type;
typedef metal::float3 type1;
struct VertexOutput {
@@ -21,11 +26,6 @@ typedef int type4;
typedef metal::float3x3 type5;
typedef metal::texturecube<float, metal::access::sample> type6;
typedef metal::sampler type7;
constexpr constant int const_2i = 2;
constexpr constant int const_1i = 1;
constexpr constant float const_4f = 4.0;
constexpr constant float const_1f = 1.0;
constexpr constant float const_0f = 0.0;
struct vs_mainInput {
};
struct vs_mainOutput {