mirror of
https://github.com/gfx-rs/wgpu.git
synced 2024-11-22 23:04:07 +00:00
[msl] inline some of the types
This commit is contained in:
Dzmitry Malyshau
Dzmitry Malyshau
parent
f1c5e04b89
commit
2774dcb403
@ -3,7 +3,7 @@ use super::{
|
||||
TranslationInfo,
|
||||
};
|
||||
use crate::{
|
||||
arena::Handle,
|
||||
arena::{Arena, Handle},
|
||||
proc::{EntryPointIndex, NameKey, Namer, TypeResolution},
|
||||
valid::{FunctionInfo, GlobalUse, ModuleInfo},
|
||||
FastHashMap,
|
||||
@ -31,6 +31,154 @@ impl Display for Level {
|
||||
}
|
||||
}
|
||||
|
||||
struct TypeContext<'a> {
|
||||
handle: Handle<crate::Type>,
|
||||
arena: &'a Arena<crate::Type>,
|
||||
names: &'a FastHashMap<NameKey, String>,
|
||||
usage: GlobalUse,
|
||||
access: crate::StorageAccess,
|
||||
first_time: bool,
|
||||
}
|
||||
|
||||
impl<'a> Display for TypeContext<'a> {
|
||||
fn fmt(&self, out: &mut Formatter<'_>) -> Result<(), FmtError> {
|
||||
let ty = &self.arena[self.handle];
|
||||
if ty.needs_alias() && !self.first_time {
|
||||
let name = &self.names[&NameKey::Type(self.handle)];
|
||||
return write!(out, "{}", name);
|
||||
}
|
||||
|
||||
match ty.inner {
|
||||
// work around Metal toolchain bug with `uint` typedef
|
||||
crate::TypeInner::Scalar {
|
||||
kind: crate::ScalarKind::Uint,
|
||||
..
|
||||
} => {
|
||||
write!(out, "metal::uint")
|
||||
}
|
||||
crate::TypeInner::Scalar { kind, .. } => {
|
||||
write!(out, "{}", scalar_kind_string(kind))
|
||||
}
|
||||
crate::TypeInner::Vector { size, kind, .. } => {
|
||||
write!(
|
||||
out,
|
||||
"{}::{}{}",
|
||||
NAMESPACE,
|
||||
scalar_kind_string(kind),
|
||||
vector_size_string(size),
|
||||
)
|
||||
}
|
||||
crate::TypeInner::Matrix { columns, rows, .. } => {
|
||||
write!(
|
||||
out,
|
||||
"{}::{}{}x{}",
|
||||
NAMESPACE,
|
||||
scalar_kind_string(crate::ScalarKind::Float),
|
||||
vector_size_string(columns),
|
||||
vector_size_string(rows),
|
||||
)
|
||||
}
|
||||
crate::TypeInner::Pointer { base, class } => {
|
||||
let sub = Self {
|
||||
arena: self.arena,
|
||||
names: self.names,
|
||||
handle: base,
|
||||
usage: self.usage,
|
||||
access: self.access,
|
||||
first_time: false,
|
||||
};
|
||||
let class_name = match class.get_name(self.usage) {
|
||||
Some(name) => name,
|
||||
None => return Ok(()),
|
||||
};
|
||||
write!(out, "{} {}*", class_name, sub)
|
||||
}
|
||||
crate::TypeInner::ValuePointer {
|
||||
size: None,
|
||||
kind,
|
||||
width: _,
|
||||
class,
|
||||
} => {
|
||||
let class_name = match class.get_name(self.usage) {
|
||||
Some(name) => name,
|
||||
None => return Ok(()),
|
||||
};
|
||||
write!(out, "{} {}*", class_name, scalar_kind_string(kind),)
|
||||
}
|
||||
crate::TypeInner::ValuePointer {
|
||||
size: Some(size),
|
||||
kind,
|
||||
width: _,
|
||||
class,
|
||||
} => {
|
||||
let class_name = match class.get_name(self.usage) {
|
||||
Some(name) => name,
|
||||
None => return Ok(()),
|
||||
};
|
||||
write!(
|
||||
out,
|
||||
"{} {}::{}{}*",
|
||||
class_name,
|
||||
NAMESPACE,
|
||||
scalar_kind_string(kind),
|
||||
vector_size_string(size),
|
||||
)
|
||||
}
|
||||
crate::TypeInner::Array { .. } | crate::TypeInner::Struct { .. } => unreachable!(),
|
||||
crate::TypeInner::Image {
|
||||
dim,
|
||||
arrayed,
|
||||
class,
|
||||
} => {
|
||||
let dim_str = match dim {
|
||||
crate::ImageDimension::D1 => "1d",
|
||||
crate::ImageDimension::D2 => "2d",
|
||||
crate::ImageDimension::D3 => "3d",
|
||||
crate::ImageDimension::Cube => "cube",
|
||||
};
|
||||
let (texture_str, msaa_str, kind, access) = match class {
|
||||
crate::ImageClass::Sampled { kind, multi } => {
|
||||
("texture", if multi { "_ms" } else { "" }, kind, "sample")
|
||||
}
|
||||
crate::ImageClass::Depth => ("depth", "", crate::ScalarKind::Float, "sample"),
|
||||
crate::ImageClass::Storage(format) => {
|
||||
let access = if self
|
||||
.access
|
||||
.contains(crate::StorageAccess::LOAD | crate::StorageAccess::STORE)
|
||||
{
|
||||
"read_write"
|
||||
} else if self.access.contains(crate::StorageAccess::STORE) {
|
||||
"write"
|
||||
} else if self.access.contains(crate::StorageAccess::LOAD) {
|
||||
"read"
|
||||
} else {
|
||||
unreachable!("module is not valid")
|
||||
};
|
||||
("texture", "", format.into(), access)
|
||||
}
|
||||
};
|
||||
let base_name = scalar_kind_string(kind);
|
||||
let array_str = if arrayed { "_array" } else { "" };
|
||||
write!(
|
||||
out,
|
||||
"{}::{}{}{}{}<{}, {}::access::{}>",
|
||||
NAMESPACE,
|
||||
texture_str,
|
||||
dim_str,
|
||||
msaa_str,
|
||||
array_str,
|
||||
base_name,
|
||||
NAMESPACE,
|
||||
access,
|
||||
)
|
||||
}
|
||||
crate::TypeInner::Sampler { comparison: _ } => {
|
||||
write!(out, "{}::sampler", NAMESPACE)
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
struct TypedGlobalVariable<'a> {
|
||||
module: &'a crate::Module,
|
||||
names: &'a FastHashMap<NameKey, String>,
|
||||
@ -43,7 +191,14 @@ impl<'a> TypedGlobalVariable<'a> {
|
||||
fn try_fmt<W: Write>(&self, out: &mut W) -> Result<(), Error> {
|
||||
let var = &self.module.global_variables[self.handle];
|
||||
let name = &self.names[&NameKey::GlobalVariable(self.handle)];
|
||||
let ty_name = &self.names[&NameKey::Type(var.ty)];
|
||||
let ty_name = TypeContext {
|
||||
handle: var.ty,
|
||||
arena: &self.module.types,
|
||||
names: self.names,
|
||||
usage: self.usage,
|
||||
access: var.storage_access,
|
||||
first_time: false,
|
||||
};
|
||||
|
||||
let (space, access, reference) = match var.class.get_name(self.usage) {
|
||||
Some(space) if self.reference => {
|
||||
@ -143,6 +298,25 @@ impl crate::StorageClass {
|
||||
}
|
||||
}
|
||||
|
||||
impl crate::Type {
|
||||
// Returns `true` if we need to emit a type alias for this type.
|
||||
fn needs_alias(&self) -> bool {
|
||||
use crate::TypeInner as Ti;
|
||||
match self.inner {
|
||||
// value types are concise enough, we only alias them if they are named
|
||||
Ti::Scalar { .. }
|
||||
| Ti::Vector { .. }
|
||||
| Ti::Matrix { .. }
|
||||
| Ti::Pointer { .. }
|
||||
| Ti::ValuePointer { .. } => self.name.is_some(),
|
||||
// composite types are better to be aliased, regardless of the name
|
||||
Ti::Struct { .. } | Ti::Array { .. } => true,
|
||||
// handle types may be different, depending on the global var access, so we always inline them
|
||||
Ti::Image { .. } | Ti::Sampler { .. } => false,
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
enum FunctionOrigin {
|
||||
Handle(Handle<crate::Function>),
|
||||
EntryPoint(EntryPointIndex),
|
||||
@ -830,7 +1004,14 @@ impl<W: Write> Writer<W> {
|
||||
) -> Result<(), Error> {
|
||||
match context.info[handle].ty {
|
||||
TypeResolution::Handle(ty_handle) => {
|
||||
let ty_name = &self.names[&NameKey::Type(ty_handle)];
|
||||
let ty_name = TypeContext {
|
||||
handle: ty_handle,
|
||||
arena: &context.module.types,
|
||||
names: &self.names,
|
||||
usage: GlobalUse::all(),
|
||||
access: crate::StorageAccess::empty(),
|
||||
first_time: false,
|
||||
};
|
||||
write!(self.out, "{}", ty_name)?;
|
||||
}
|
||||
TypeResolution::Value(crate::TypeInner::Scalar { kind, .. }) => {
|
||||
@ -1096,92 +1277,25 @@ impl<W: Write> Writer<W> {
|
||||
|
||||
fn write_type_defs(&mut self, module: &crate::Module) -> Result<(), Error> {
|
||||
for (handle, ty) in module.types.iter() {
|
||||
if !ty.needs_alias() {
|
||||
continue;
|
||||
}
|
||||
let name = &self.names[&NameKey::Type(handle)];
|
||||
let global_use = GlobalUse::all(); //TODO
|
||||
match ty.inner {
|
||||
// work around Metal toolchain bug with `uint` typedef
|
||||
crate::TypeInner::Scalar {
|
||||
kind: crate::ScalarKind::Uint,
|
||||
..
|
||||
} => {
|
||||
writeln!(self.out, "typedef metal::uint {};", name)?;
|
||||
}
|
||||
crate::TypeInner::Scalar { kind, .. } => {
|
||||
writeln!(self.out, "typedef {} {};", scalar_kind_string(kind), name)?;
|
||||
}
|
||||
crate::TypeInner::Vector { size, kind, .. } => {
|
||||
writeln!(
|
||||
self.out,
|
||||
"typedef {}::{}{} {};",
|
||||
NAMESPACE,
|
||||
scalar_kind_string(kind),
|
||||
vector_size_string(size),
|
||||
name
|
||||
)?;
|
||||
}
|
||||
crate::TypeInner::Matrix { columns, rows, .. } => {
|
||||
writeln!(
|
||||
self.out,
|
||||
"typedef {}::{}{}x{} {};",
|
||||
NAMESPACE,
|
||||
scalar_kind_string(crate::ScalarKind::Float),
|
||||
vector_size_string(columns),
|
||||
vector_size_string(rows),
|
||||
name
|
||||
)?;
|
||||
}
|
||||
crate::TypeInner::Pointer { base, class } => {
|
||||
let base_name = &self.names[&NameKey::Type(base)];
|
||||
let class_name = match class.get_name(global_use) {
|
||||
Some(name) => name,
|
||||
None => continue,
|
||||
};
|
||||
writeln!(self.out, "typedef {} {} *{};", class_name, base_name, name)?;
|
||||
}
|
||||
crate::TypeInner::ValuePointer {
|
||||
size: None,
|
||||
kind,
|
||||
width: _,
|
||||
class,
|
||||
} => {
|
||||
let class_name = match class.get_name(global_use) {
|
||||
Some(name) => name,
|
||||
None => continue,
|
||||
};
|
||||
writeln!(
|
||||
self.out,
|
||||
"typedef {} {} *{};",
|
||||
class_name,
|
||||
scalar_kind_string(kind),
|
||||
name
|
||||
)?;
|
||||
}
|
||||
crate::TypeInner::ValuePointer {
|
||||
size: Some(size),
|
||||
kind,
|
||||
width: _,
|
||||
class,
|
||||
} => {
|
||||
let class_name = match class.get_name(global_use) {
|
||||
Some(name) => name,
|
||||
None => continue,
|
||||
};
|
||||
writeln!(
|
||||
self.out,
|
||||
"typedef {} {}::{}{} {};",
|
||||
class_name,
|
||||
NAMESPACE,
|
||||
scalar_kind_string(kind),
|
||||
vector_size_string(size),
|
||||
name
|
||||
)?;
|
||||
}
|
||||
crate::TypeInner::Array {
|
||||
base,
|
||||
size,
|
||||
stride: _,
|
||||
} => {
|
||||
let base_name = &self.names[&NameKey::Type(base)];
|
||||
let base_name = TypeContext {
|
||||
handle: base,
|
||||
arena: &module.types,
|
||||
names: &self.names,
|
||||
usage: global_use,
|
||||
access: crate::StorageAccess::empty(),
|
||||
first_time: false,
|
||||
};
|
||||
let size_str = match size {
|
||||
crate::ArraySize::Constant(const_handle) => {
|
||||
&self.names[&NameKey::Constant(const_handle)]
|
||||
@ -1197,68 +1311,28 @@ impl<W: Write> Writer<W> {
|
||||
writeln!(self.out, "struct {} {{", name)?;
|
||||
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)];
|
||||
let base_name = TypeContext {
|
||||
handle: member.ty,
|
||||
arena: &module.types,
|
||||
names: &self.names,
|
||||
usage: global_use,
|
||||
access: crate::StorageAccess::empty(),
|
||||
first_time: false,
|
||||
};
|
||||
writeln!(self.out, "{}{} {};", INDENT, base_name, member_name)?;
|
||||
}
|
||||
writeln!(self.out, "}};")?;
|
||||
}
|
||||
crate::TypeInner::Image {
|
||||
dim,
|
||||
arrayed,
|
||||
class,
|
||||
} => {
|
||||
let dim_str = match dim {
|
||||
crate::ImageDimension::D1 => "1d",
|
||||
crate::ImageDimension::D2 => "2d",
|
||||
crate::ImageDimension::D3 => "3d",
|
||||
crate::ImageDimension::Cube => "cube",
|
||||
_ => {
|
||||
let ty_name = TypeContext {
|
||||
handle,
|
||||
arena: &module.types,
|
||||
names: &self.names,
|
||||
usage: global_use,
|
||||
access: crate::StorageAccess::empty(),
|
||||
first_time: true,
|
||||
};
|
||||
let (texture_str, msaa_str, kind, access) = match class {
|
||||
crate::ImageClass::Sampled { kind, multi } => {
|
||||
("texture", if multi { "_ms" } else { "" }, kind, "sample")
|
||||
}
|
||||
crate::ImageClass::Depth => {
|
||||
("depth", "", crate::ScalarKind::Float, "sample")
|
||||
}
|
||||
crate::ImageClass::Storage(format) => {
|
||||
let (_, global) = module
|
||||
.global_variables
|
||||
.iter()
|
||||
.find(|&(_, ref var)| var.ty == handle)
|
||||
.expect("Unable to find a global variable using the image type");
|
||||
let access = if global
|
||||
.storage_access
|
||||
.contains(crate::StorageAccess::LOAD | crate::StorageAccess::STORE)
|
||||
{
|
||||
"read_write"
|
||||
} else if global.storage_access.contains(crate::StorageAccess::STORE) {
|
||||
"write"
|
||||
} else if global.storage_access.contains(crate::StorageAccess::LOAD) {
|
||||
"read"
|
||||
} else {
|
||||
return Err(Error::Validation);
|
||||
};
|
||||
("texture", "", format.into(), access)
|
||||
}
|
||||
};
|
||||
let base_name = scalar_kind_string(kind);
|
||||
let array_str = if arrayed { "_array" } else { "" };
|
||||
writeln!(
|
||||
self.out,
|
||||
"typedef {}::{}{}{}{}<{}, {}::access::{}> {};",
|
||||
NAMESPACE,
|
||||
texture_str,
|
||||
dim_str,
|
||||
msaa_str,
|
||||
array_str,
|
||||
base_name,
|
||||
NAMESPACE,
|
||||
access,
|
||||
name
|
||||
)?;
|
||||
}
|
||||
crate::TypeInner::Sampler { comparison: _ } => {
|
||||
writeln!(self.out, "typedef {}::sampler {};", NAMESPACE, name)?;
|
||||
writeln!(self.out, "typedef {} {};", ty_name, name)?;
|
||||
}
|
||||
}
|
||||
}
|
||||
@ -1305,7 +1379,14 @@ impl<W: Write> Writer<W> {
|
||||
crate::ConstantInner::Scalar { .. } => {}
|
||||
crate::ConstantInner::Composite { ty, ref components } => {
|
||||
let name = &self.names[&NameKey::Constant(handle)];
|
||||
let ty_name = &self.names[&NameKey::Type(ty)];
|
||||
let ty_name = TypeContext {
|
||||
handle: ty,
|
||||
arena: &module.types,
|
||||
names: &self.names,
|
||||
usage: GlobalUse::empty(),
|
||||
access: crate::StorageAccess::empty(),
|
||||
first_time: false,
|
||||
};
|
||||
write!(self.out, "constexpr constant {} {} = {{", ty_name, name,)?;
|
||||
for (i, &sub_handle) in components.iter().enumerate() {
|
||||
let separator = if i != 0 { ", " } else { "" };
|
||||
@ -1415,16 +1496,36 @@ impl<W: Write> Writer<W> {
|
||||
}
|
||||
}
|
||||
|
||||
writeln!(self.out)?;
|
||||
let fun_name = &self.names[&NameKey::Function(fun_handle)];
|
||||
let result_type_name = match fun.result {
|
||||
Some(ref result) => &self.names[&NameKey::Type(result.ty)],
|
||||
None => "void",
|
||||
};
|
||||
writeln!(self.out, "{} {}(", result_type_name, fun_name)?;
|
||||
match fun.result {
|
||||
Some(ref result) => {
|
||||
let ty_name = TypeContext {
|
||||
handle: result.ty,
|
||||
arena: &module.types,
|
||||
names: &self.names,
|
||||
usage: GlobalUse::empty(),
|
||||
access: crate::StorageAccess::empty(),
|
||||
first_time: false,
|
||||
};
|
||||
write!(self.out, "{}", ty_name)?;
|
||||
}
|
||||
None => {
|
||||
write!(self.out, "void")?;
|
||||
}
|
||||
}
|
||||
writeln!(self.out, " {}(", fun_name)?;
|
||||
|
||||
for (index, arg) in fun.arguments.iter().enumerate() {
|
||||
let name = &self.names[&NameKey::FunctionArgument(fun_handle, index as u32)];
|
||||
let param_type_name = &self.names[&NameKey::Type(arg.ty)];
|
||||
let param_type_name = TypeContext {
|
||||
handle: arg.ty,
|
||||
arena: &module.types,
|
||||
names: &self.names,
|
||||
usage: GlobalUse::empty(),
|
||||
access: crate::StorageAccess::empty(),
|
||||
first_time: false,
|
||||
};
|
||||
let separator =
|
||||
separate(!pass_through_globals.is_empty() || index + 1 != fun.arguments.len());
|
||||
writeln!(
|
||||
@ -1449,7 +1550,14 @@ impl<W: Write> Writer<W> {
|
||||
writeln!(self.out, ") {{")?;
|
||||
|
||||
for (local_handle, local) in fun.local_variables.iter() {
|
||||
let ty_name = &self.names[&NameKey::Type(local.ty)];
|
||||
let ty_name = TypeContext {
|
||||
handle: local.ty,
|
||||
arena: &module.types,
|
||||
names: &self.names,
|
||||
usage: GlobalUse::empty(),
|
||||
access: crate::StorageAccess::empty(),
|
||||
first_time: false,
|
||||
};
|
||||
let local_name = &self.names[&NameKey::FunctionLocal(fun_handle, local_handle)];
|
||||
write!(self.out, "{}{} {}", INDENT, ty_name, local_name)?;
|
||||
if let Some(value) = local.init {
|
||||
@ -1473,7 +1581,6 @@ impl<W: Write> Writer<W> {
|
||||
self.named_expressions.clear();
|
||||
self.put_block(Level(1), &fun.body, &context)?;
|
||||
writeln!(self.out, "}}")?;
|
||||
writeln!(self.out)?;
|
||||
}
|
||||
|
||||
let mut info = TranslationInfo {
|
||||
@ -1503,6 +1610,7 @@ impl<W: Write> Writer<W> {
|
||||
}
|
||||
}
|
||||
|
||||
writeln!(self.out)?;
|
||||
let fun_name = &self.names[&NameKey::EntryPoint(ep_index as _)];
|
||||
info.entry_point_names.push(Ok(fun_name.clone()));
|
||||
|
||||
@ -1558,9 +1666,16 @@ impl<W: Write> Writer<W> {
|
||||
};
|
||||
varying_count += 1;
|
||||
let name = &self.names[&name_key];
|
||||
let type_name = &self.names[&NameKey::Type(ty)];
|
||||
let ty_name = TypeContext {
|
||||
handle: ty,
|
||||
arena: &module.types,
|
||||
names: &self.names,
|
||||
usage: GlobalUse::empty(),
|
||||
access: crate::StorageAccess::empty(),
|
||||
first_time: false,
|
||||
};
|
||||
let resolved = options.resolve_local_binding(binding, in_mode)?;
|
||||
write!(self.out, "{}{} {}", INDENT, type_name, name)?;
|
||||
write!(self.out, "{}{} {}", INDENT, ty_name, name)?;
|
||||
resolved.try_fmt_decorated(&mut self.out, "")?;
|
||||
writeln!(self.out, ";")?;
|
||||
}
|
||||
@ -1592,7 +1707,14 @@ impl<W: Write> Writer<W> {
|
||||
}
|
||||
writeln!(self.out, "struct {} {{", stage_out_name)?;
|
||||
for (name, ty, binding) in result_members {
|
||||
let type_name = &self.names[&NameKey::Type(ty)];
|
||||
let ty_name = TypeContext {
|
||||
handle: ty,
|
||||
arena: &module.types,
|
||||
names: &self.names,
|
||||
usage: GlobalUse::empty(),
|
||||
access: crate::StorageAccess::empty(),
|
||||
first_time: false,
|
||||
};
|
||||
let binding = binding.ok_or(Error::Validation)?;
|
||||
if !pipeline_options.allow_point_size
|
||||
&& *binding == crate::Binding::BuiltIn(crate::BuiltIn::PointSize)
|
||||
@ -1600,7 +1722,7 @@ impl<W: Write> Writer<W> {
|
||||
continue;
|
||||
}
|
||||
let resolved = options.resolve_local_binding(binding, out_mode)?;
|
||||
write!(self.out, "{}{} {}", INDENT, type_name, name)?;
|
||||
write!(self.out, "{}{} {}", INDENT, ty_name, name)?;
|
||||
resolved.try_fmt_decorated(&mut self.out, "")?;
|
||||
writeln!(self.out, ";")?;
|
||||
}
|
||||
@ -1626,7 +1748,14 @@ impl<W: Write> Writer<W> {
|
||||
_ => continue,
|
||||
};
|
||||
let name = &self.names[&name_key];
|
||||
let type_name = &self.names[&NameKey::Type(ty)];
|
||||
let ty_name = TypeContext {
|
||||
handle: ty,
|
||||
arena: &module.types,
|
||||
names: &self.names,
|
||||
usage: GlobalUse::empty(),
|
||||
access: crate::StorageAccess::empty(),
|
||||
first_time: false,
|
||||
};
|
||||
let resolved = options.resolve_local_binding(binding, in_mode)?;
|
||||
let separator = if is_first_argument {
|
||||
is_first_argument = false;
|
||||
@ -1634,7 +1763,7 @@ impl<W: Write> Writer<W> {
|
||||
} else {
|
||||
','
|
||||
};
|
||||
write!(self.out, "{} {} {}", separator, type_name, name)?;
|
||||
write!(self.out, "{} {} {}", separator, ty_name, name)?;
|
||||
resolved.try_fmt_decorated(&mut self.out, "\n")?;
|
||||
}
|
||||
for (handle, var) in module.global_variables.iter() {
|
||||
@ -1763,7 +1892,14 @@ impl<W: Write> Writer<W> {
|
||||
//TODO: we can postpone this till the relevant expressions are emitted
|
||||
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)];
|
||||
let ty_name = TypeContext {
|
||||
handle: local.ty,
|
||||
arena: &module.types,
|
||||
names: &self.names,
|
||||
usage: GlobalUse::empty(),
|
||||
access: crate::StorageAccess::empty(),
|
||||
first_time: false,
|
||||
};
|
||||
write!(self.out, "{}{} {}", INDENT, ty_name, name)?;
|
||||
if let Some(value) = local.init {
|
||||
let value_str = &self.names[&NameKey::Constant(value)];
|
||||
|
||||
@ -14,46 +14,42 @@ 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 metal::uint type;
|
||||
typedef metal::float2 type1;
|
||||
struct Particle {
|
||||
type1 pos;
|
||||
type1 vel;
|
||||
metal::float2 pos;
|
||||
metal::float2 vel;
|
||||
};
|
||||
typedef float type2;
|
||||
struct SimParams {
|
||||
type2 deltaT;
|
||||
type2 rule1Distance;
|
||||
type2 rule2Distance;
|
||||
type2 rule3Distance;
|
||||
type2 rule1Scale;
|
||||
type2 rule2Scale;
|
||||
type2 rule3Scale;
|
||||
float deltaT;
|
||||
float rule1Distance;
|
||||
float rule2Distance;
|
||||
float rule3Distance;
|
||||
float rule1Scale;
|
||||
float rule2Scale;
|
||||
float rule3Scale;
|
||||
};
|
||||
typedef Particle type3[1];
|
||||
struct Particles {
|
||||
type3 particles;
|
||||
};
|
||||
typedef metal::uint3 type4;
|
||||
typedef int type5;
|
||||
|
||||
struct main1Input {
|
||||
};
|
||||
kernel void main1(
|
||||
type4 global_invocation_id [[thread_position_in_grid]]
|
||||
metal::uint3 global_invocation_id [[thread_position_in_grid]]
|
||||
, constant SimParams& params [[buffer(0)]]
|
||||
, constant Particles& particlesSrc [[buffer(1)]]
|
||||
, device Particles& particlesDst [[buffer(2)]]
|
||||
) {
|
||||
type1 vPos;
|
||||
type1 vVel;
|
||||
type1 cMass;
|
||||
type1 cVel;
|
||||
type1 colVel;
|
||||
type5 cMassCount = const_0i;
|
||||
type5 cVelCount = const_0i;
|
||||
type1 pos1;
|
||||
type1 vel1;
|
||||
type i = const_0u;
|
||||
metal::float2 vPos;
|
||||
metal::float2 vVel;
|
||||
metal::float2 cMass;
|
||||
metal::float2 cVel;
|
||||
metal::float2 colVel;
|
||||
int cMassCount = const_0i;
|
||||
int cVelCount = const_0i;
|
||||
metal::float2 pos1;
|
||||
metal::float2 vel1;
|
||||
metal::uint i = const_0u;
|
||||
if (global_invocation_id.x >= NUM_PARTICLES) {
|
||||
return;
|
||||
}
|
||||
|
||||
@ -9,17 +9,16 @@ 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 metal::uint type;
|
||||
typedef type type1[1];
|
||||
typedef metal::uint type1[1];
|
||||
struct PrimeIndices {
|
||||
type1 data;
|
||||
};
|
||||
typedef metal::uint3 type2;
|
||||
type collatz_iterations(
|
||||
type n_base
|
||||
|
||||
metal::uint collatz_iterations(
|
||||
metal::uint n_base
|
||||
) {
|
||||
type n;
|
||||
type i = const_0u;
|
||||
metal::uint n;
|
||||
metal::uint i = const_0u;
|
||||
n = n_base;
|
||||
while(true) {
|
||||
if (n <= const_1u) {
|
||||
@ -38,10 +37,10 @@ type collatz_iterations(
|
||||
struct main1Input {
|
||||
};
|
||||
kernel void main1(
|
||||
type2 global_id [[thread_position_in_grid]]
|
||||
metal::uint3 global_id [[thread_position_in_grid]]
|
||||
, device PrimeIndices& v_indices [[user(fake0)]]
|
||||
) {
|
||||
type _expr9 = collatz_iterations(v_indices.data[global_id.x]);
|
||||
metal::uint _expr9 = collatz_iterations(v_indices.data[global_id.x]);
|
||||
v_indices.data[global_id.x] = _expr9;
|
||||
return;
|
||||
}
|
||||
|
||||
@ -5,6 +5,7 @@ expression: msl
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
|
||||
kernel void main1(
|
||||
) {
|
||||
return;
|
||||
|
||||
@ -7,17 +7,13 @@ expression: msl
|
||||
|
||||
constexpr constant int const_10i = 10;
|
||||
constexpr constant int const_20i = 20;
|
||||
typedef metal::texture2d<uint, metal::access::read> type;
|
||||
typedef metal::texture1d<uint, metal::access::write> type1;
|
||||
typedef metal::uint3 type2;
|
||||
typedef metal::uint2 type3;
|
||||
typedef metal::int2 type4;
|
||||
|
||||
struct main1Input {
|
||||
};
|
||||
kernel void main1(
|
||||
type2 local_id [[thread_position_in_threadgroup]]
|
||||
, type image_src [[user(fake0)]]
|
||||
, type1 image_dst [[user(fake0)]]
|
||||
metal::uint3 local_id [[thread_position_in_threadgroup]]
|
||||
, metal::texture2d<uint, metal::access::read> image_src [[user(fake0)]]
|
||||
, metal::texture1d<uint, metal::access::write> image_dst [[user(fake0)]]
|
||||
) {
|
||||
metal::int2 _expr12 = (int2(image_src.get_width(), image_src.get_height()) * static_cast<int2>(metal::uint2(local_id.x, local_id.y))) % metal::int2(const_10i, const_20i);
|
||||
metal::uint4 _expr13 = image_src.read(metal::uint2(_expr12));
|
||||
|
||||
@ -13,57 +13,49 @@ constexpr constant unsigned const_1u = 1u;
|
||||
constexpr constant int const_0i1 = 0;
|
||||
constexpr constant float const_0f = 0.0;
|
||||
constexpr constant float const_1f = 1.0;
|
||||
typedef float type;
|
||||
typedef metal::float2 type1;
|
||||
typedef thread type1 *type2;
|
||||
typedef thread type1 *type3;
|
||||
typedef metal::float4 type4;
|
||||
typedef metal::uint type5;
|
||||
typedef type type6[const_1u];
|
||||
typedef float type6[const_1u];
|
||||
struct gl_PerVertex {
|
||||
type4 gl_Position;
|
||||
type gl_PointSize;
|
||||
metal::float4 gl_Position;
|
||||
float gl_PointSize;
|
||||
type6 gl_ClipDistance;
|
||||
type6 gl_CullDistance;
|
||||
};
|
||||
typedef thread gl_PerVertex *type7;
|
||||
typedef int type8;
|
||||
typedef thread type4 *type9;
|
||||
struct type10 {
|
||||
type1 member;
|
||||
type4 gl_Position1;
|
||||
type gl_PointSize1;
|
||||
metal::float2 member;
|
||||
metal::float4 gl_Position1;
|
||||
float gl_PointSize1;
|
||||
type6 gl_ClipDistance1;
|
||||
};
|
||||
|
||||
void main1(
|
||||
thread type1& v_uv,
|
||||
thread type1 const& a_uv,
|
||||
thread metal::float2& v_uv,
|
||||
thread metal::float2 const& a_uv,
|
||||
thread gl_PerVertex& _,
|
||||
thread type1 const& a_pos
|
||||
thread metal::float2 const& a_pos
|
||||
) {
|
||||
v_uv = a_uv;
|
||||
type1 _expr13 = a_pos;
|
||||
metal::float2 _expr13 = a_pos;
|
||||
_.gl_Position = metal::float4(_expr13.x, _expr13.y, const_0f, const_1f);
|
||||
return;
|
||||
}
|
||||
|
||||
struct main2Input {
|
||||
type1 a_uv1 [[attribute(1)]];
|
||||
type1 a_pos1 [[attribute(0)]];
|
||||
metal::float2 a_uv1 [[attribute(1)]];
|
||||
metal::float2 a_pos1 [[attribute(0)]];
|
||||
};
|
||||
struct main2Output {
|
||||
type1 member [[user(loc0)]];
|
||||
type4 gl_Position1 [[position]];
|
||||
type gl_PointSize1 [[point_size]];
|
||||
metal::float2 member [[user(loc0)]];
|
||||
metal::float4 gl_Position1 [[position]];
|
||||
float gl_PointSize1 [[point_size]];
|
||||
type6 gl_ClipDistance1 [[clip_distance]];
|
||||
};
|
||||
vertex main2Output main2(
|
||||
main2Input varyings [[stage_in]]
|
||||
) {
|
||||
type1 v_uv = {};
|
||||
type1 a_uv = {};
|
||||
metal::float2 v_uv = {};
|
||||
metal::float2 a_uv = {};
|
||||
gl_PerVertex _ = {};
|
||||
type1 a_pos = {};
|
||||
metal::float2 a_pos = {};
|
||||
const auto a_uv1 = varyings.a_uv1;
|
||||
const auto a_pos1 = varyings.a_pos1;
|
||||
a_uv = a_uv1;
|
||||
|
||||
@ -8,22 +8,18 @@ expression: msl
|
||||
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;
|
||||
struct VertexOutput {
|
||||
type1 uv;
|
||||
type2 position;
|
||||
metal::float2 uv;
|
||||
metal::float4 position;
|
||||
};
|
||||
typedef metal::texture2d<float, metal::access::sample> type3;
|
||||
typedef metal::sampler type4;
|
||||
|
||||
struct main1Input {
|
||||
type1 pos [[attribute(0)]];
|
||||
type1 uv1 [[attribute(1)]];
|
||||
metal::float2 pos [[attribute(0)]];
|
||||
metal::float2 uv1 [[attribute(1)]];
|
||||
};
|
||||
struct main1Output {
|
||||
type1 uv [[user(loc0)]];
|
||||
type2 position [[position]];
|
||||
metal::float2 uv [[user(loc0)]];
|
||||
metal::float4 position [[position]];
|
||||
};
|
||||
vertex main1Output main1(
|
||||
main1Input varyings [[stage_in]]
|
||||
@ -37,16 +33,17 @@ vertex main1Output main1(
|
||||
return main1Output { _tmp.uv, _tmp.position };
|
||||
}
|
||||
|
||||
|
||||
struct main2Input {
|
||||
type1 uv2 [[user(loc0)]];
|
||||
metal::float2 uv2 [[user(loc0)]];
|
||||
};
|
||||
struct main2Output {
|
||||
type2 member1 [[color(0)]];
|
||||
metal::float4 member1 [[color(0)]];
|
||||
};
|
||||
fragment main2Output main2(
|
||||
main2Input varyings1 [[stage_in]]
|
||||
, type3 u_texture [[user(fake0)]]
|
||||
, type4 u_sampler [[user(fake0)]]
|
||||
, metal::texture2d<float, metal::access::sample> u_texture [[user(fake0)]]
|
||||
, metal::sampler u_sampler [[user(fake0)]]
|
||||
) {
|
||||
const auto uv2 = varyings1.uv2;
|
||||
metal::float4 _expr4 = u_texture.sample(u_sampler, uv2);
|
||||
|
||||
@ -13,33 +13,25 @@ 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;
|
||||
metal::uint4 num_lights;
|
||||
};
|
||||
typedef metal::float4x4 type1;
|
||||
typedef metal::float4 type2;
|
||||
struct Light {
|
||||
type1 proj;
|
||||
type2 pos;
|
||||
type2 color;
|
||||
metal::float4x4 proj;
|
||||
metal::float4 pos;
|
||||
metal::float4 color;
|
||||
};
|
||||
typedef Light type3[1];
|
||||
struct Lights {
|
||||
type3 data;
|
||||
};
|
||||
typedef metal::depth2d_array<float, metal::access::sample> type4;
|
||||
typedef metal::sampler type5;
|
||||
typedef metal::uint type6;
|
||||
typedef float type7;
|
||||
typedef metal::float2 type8;
|
||||
typedef metal::float3 type9;
|
||||
constexpr constant type9 c_ambient = {const_0_05f, const_0_05f, const_0_05f};
|
||||
type7 fetch_shadow(
|
||||
type6 light_id,
|
||||
type2 homogeneous_coords,
|
||||
type4 t_shadow,
|
||||
type5 sampler_shadow
|
||||
constexpr constant metal::float3 c_ambient = {const_0_05f, const_0_05f, const_0_05f};
|
||||
|
||||
float fetch_shadow(
|
||||
metal::uint light_id,
|
||||
metal::float4 homogeneous_coords,
|
||||
metal::depth2d_array<float, metal::access::sample> t_shadow,
|
||||
metal::sampler sampler_shadow
|
||||
) {
|
||||
if (homogeneous_coords.w <= const_0f) {
|
||||
return const_1f;
|
||||
@ -50,23 +42,23 @@ type7 fetch_shadow(
|
||||
}
|
||||
|
||||
struct fs_mainInput {
|
||||
type9 raw_normal [[user(loc0)]];
|
||||
type2 position [[user(loc1)]];
|
||||
metal::float3 raw_normal [[user(loc0)]];
|
||||
metal::float4 position [[user(loc1)]];
|
||||
};
|
||||
struct fs_mainOutput {
|
||||
type2 member [[color(0)]];
|
||||
metal::float4 member [[color(0)]];
|
||||
};
|
||||
fragment fs_mainOutput fs_main(
|
||||
fs_mainInput varyings [[stage_in]]
|
||||
, constant Globals& u_globals [[user(fake0)]]
|
||||
, constant Lights& s_lights [[user(fake0)]]
|
||||
, type4 t_shadow [[user(fake0)]]
|
||||
, type5 sampler_shadow [[user(fake0)]]
|
||||
, metal::depth2d_array<float, metal::access::sample> t_shadow [[user(fake0)]]
|
||||
, metal::sampler sampler_shadow [[user(fake0)]]
|
||||
) {
|
||||
const auto raw_normal = varyings.raw_normal;
|
||||
const auto position = varyings.position;
|
||||
type9 color1 = c_ambient;
|
||||
type6 i = const_0u;
|
||||
metal::float3 color1 = c_ambient;
|
||||
metal::uint i = const_0u;
|
||||
bool loop_init = true;
|
||||
while(true) {
|
||||
if (!loop_init) {
|
||||
@ -77,7 +69,7 @@ fragment fs_mainOutput fs_main(
|
||||
break;
|
||||
}
|
||||
Light _expr21 = s_lights.data[i];
|
||||
type7 _expr25 = fetch_shadow(i, _expr21.proj * position, t_shadow, sampler_shadow);
|
||||
float _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) };
|
||||
|
||||
@ -10,38 +10,31 @@ 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 {
|
||||
type position;
|
||||
type1 uv;
|
||||
metal::float4 position;
|
||||
metal::float3 uv;
|
||||
};
|
||||
typedef metal::float4x4 type2;
|
||||
struct Data {
|
||||
type2 proj_inv;
|
||||
type2 view;
|
||||
metal::float4x4 proj_inv;
|
||||
metal::float4x4 view;
|
||||
};
|
||||
typedef metal::uint type3;
|
||||
typedef int type4;
|
||||
typedef metal::float3x3 type5;
|
||||
typedef metal::texturecube<float, metal::access::sample> type6;
|
||||
typedef metal::sampler type7;
|
||||
|
||||
struct vs_mainInput {
|
||||
};
|
||||
struct vs_mainOutput {
|
||||
type position [[position]];
|
||||
type1 uv [[user(loc0)]];
|
||||
metal::float4 position [[position]];
|
||||
metal::float3 uv [[user(loc0)]];
|
||||
};
|
||||
vertex vs_mainOutput vs_main(
|
||||
type3 vertex_index [[vertex_id]]
|
||||
metal::uint vertex_index [[vertex_id]]
|
||||
, constant Data& r_data [[buffer(0)]]
|
||||
) {
|
||||
type4 tmp1_;
|
||||
type4 tmp2_;
|
||||
int tmp1_;
|
||||
int 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 _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;
|
||||
@ -49,16 +42,17 @@ vertex vs_mainOutput vs_main(
|
||||
return vs_mainOutput { _tmp.position, _tmp.uv };
|
||||
}
|
||||
|
||||
|
||||
struct fs_mainInput {
|
||||
type1 uv [[user(loc0)]];
|
||||
metal::float3 uv [[user(loc0)]];
|
||||
};
|
||||
struct fs_mainOutput {
|
||||
type member1 [[color(0)]];
|
||||
metal::float4 member1 [[color(0)]];
|
||||
};
|
||||
fragment fs_mainOutput fs_main(
|
||||
fs_mainInput varyings1 [[stage_in]]
|
||||
, type position [[position]]
|
||||
, type6 r_texture [[texture(0)]]
|
||||
, metal::float4 position [[position]]
|
||||
, metal::texturecube<float, metal::access::sample> r_texture [[texture(0)]]
|
||||
) {
|
||||
constexpr metal::sampler r_sampler(
|
||||
metal::s_address::clamp_to_edge,
|
||||
|
||||
Loading…
Reference in New Issue
Block a user