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

[wgsl-in] Handle modf and frexp (#2454)

Browse Source
This commit is contained in:
Fredrik Fornwall
2023-09-02 08:55:35 +02:00
committed by GitHub
parent f49314dbbd
commit 5329aa2bd5
26 changed files with 843 additions and 205 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
src/back/glsl/keywords.rs
View File

@@ -477,4 +477,7 @@ pub const RESERVED_KEYWORDS: &[&str] = &[
// entry point name (should not be shadowed)
//
"main",
// Naga utilities:
super::MODF_FUNCTION,
super::FREXP_FUNCTION,
];

54
src/back/glsl/mod.rs
View File

@@ -72,6 +72,9 @@ pub const SUPPORTED_ES_VERSIONS: &[u16] = &[300, 310, 320];
/// of detail for bounds checking in `ImageLoad`
const CLAMPED_LOD_SUFFIX: &str = "_clamped_lod";
pub(crate) const MODF_FUNCTION: &str = "naga_modf";
pub(crate) const FREXP_FUNCTION: &str = "naga_frexp";
/// Mapping between resources and bindings.
pub type BindingMap = std::collections::BTreeMap<crate::ResourceBinding, u8>;
@@ -631,6 +634,53 @@ impl<'a, W: Write> Writer<'a, W> {
}
}
// Write functions to create special types.
for (type_key, struct_ty) in self.module.special_types.predeclared_types.iter() {
match type_key {
&crate::PredeclaredType::ModfResult { size, width }
| &crate::PredeclaredType::FrexpResult { size, width } => {
let arg_type_name_owner;
let arg_type_name = if let Some(size) = size {
arg_type_name_owner =
format!("{}vec{}", if width == 8 { "d" } else { "" }, size as u8);
&arg_type_name_owner
} else if width == 8 {
"double"
} else {
"float"
};
let other_type_name_owner;
let (defined_func_name, called_func_name, other_type_name) =
if matches!(type_key, &crate::PredeclaredType::ModfResult { .. }) {
(MODF_FUNCTION, "modf", arg_type_name)
} else {
let other_type_name = if let Some(size) = size {
other_type_name_owner = format!("ivec{}", size as u8);
&other_type_name_owner
} else {
"int"
};
(FREXP_FUNCTION, "frexp", other_type_name)
};
let struct_name = &self.names[&NameKey::Type(*struct_ty)];
writeln!(self.out)?;
writeln!(
self.out,
"{} {defined_func_name}({arg_type_name} arg) {{
{other_type_name} other;
{arg_type_name} fract = {called_func_name}(arg, other);
return {}(fract, other);
}}",
struct_name, struct_name
)?;
}
&crate::PredeclaredType::AtomicCompareExchangeWeakResult { .. } => {}
}
}
// Write all named constants
let mut constants = self
.module
@@ -2997,8 +3047,8 @@ impl<'a, W: Write> Writer<'a, W> {
Mf::Round => "roundEven",
Mf::Fract => "fract",
Mf::Trunc => "trunc",
Mf::Modf => "modf",
Mf::Frexp => "frexp",
Mf::Modf => MODF_FUNCTION,
Mf::Frexp => FREXP_FUNCTION,
Mf::Ldexp => "ldexp",
// exponent
Mf::Exp => "exp",

53
src/back/hlsl/help.rs
View File

@@ -781,6 +781,59 @@ impl<'a, W: Write> super::Writer<'a, W> {
Ok(())
}
/// Write functions to create special types.
pub(super) fn write_special_functions(&mut self, module: &crate::Module) -> BackendResult {
for (type_key, struct_ty) in module.special_types.predeclared_types.iter() {
match type_key {
&crate::PredeclaredType::ModfResult { size, width }
| &crate::PredeclaredType::FrexpResult { size, width } => {
let arg_type_name_owner;
let arg_type_name = if let Some(size) = size {
arg_type_name_owner = format!(
"{}{}",
if width == 8 { "double" } else { "float" },
size as u8
);
&arg_type_name_owner
} else if width == 8 {
"double"
} else {
"float"
};
let (defined_func_name, called_func_name, second_field_name, sign_multiplier) =
if matches!(type_key, &crate::PredeclaredType::ModfResult { .. }) {
(super::writer::MODF_FUNCTION, "modf", "whole", "")
} else {
(
super::writer::FREXP_FUNCTION,
"frexp",
"exp_",
"sign(arg) * ",
)
};
let struct_name = &self.names[&NameKey::Type(*struct_ty)];
writeln!(
self.out,
"{struct_name} {defined_func_name}({arg_type_name} arg) {{
{arg_type_name} other;
{struct_name} result;
result.fract = {sign_multiplier}{called_func_name}(arg, other);
result.{second_field_name} = other;
return result;
}}"
)?;
writeln!(self.out)?;
}
&crate::PredeclaredType::AtomicCompareExchangeWeakResult { .. } => {}
}
}
Ok(())
}
/// Helper function that writes compose wrapped functions
pub(super) fn write_wrapped_compose_functions(
&mut self,

3
src/back/hlsl/keywords.rs
View File

@@ -814,6 +814,9 @@ pub const RESERVED: &[&str] = &[
"TextureBuffer",
"ConstantBuffer",
"RayQuery",
// Naga utilities
super::writer::MODF_FUNCTION,
super::writer::FREXP_FUNCTION,
];
// DXC scalar types, from https://github.com/microsoft/DirectXShaderCompiler/blob/18c9e114f9c314f93e68fbc72ce207d4ed2e65ae/tools/clang/lib/AST/ASTContextHLSL.cpp#L48-L254

9
src/back/hlsl/writer.rs
View File

@@ -17,6 +17,9 @@ const SPECIAL_BASE_VERTEX: &str = "base_vertex";
const SPECIAL_BASE_INSTANCE: &str = "base_instance";
const SPECIAL_OTHER: &str = "other";
pub(crate) const MODF_FUNCTION: &str = "naga_modf";
pub(crate) const FREXP_FUNCTION: &str = "naga_frexp";
struct EpStructMember {
name: String,
ty: Handle<crate::Type>,
@@ -244,6 +247,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
}
}
self.write_special_functions(module)?;
self.write_wrapped_compose_functions(module, &module.const_expressions)?;
// Write all named constants
@@ -2675,8 +2680,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
Mf::Round => Function::Regular("round"),
Mf::Fract => Function::Regular("frac"),
Mf::Trunc => Function::Regular("trunc"),
Mf::Modf => Function::Regular("modf"),
Mf::Frexp => Function::Regular("frexp"),
Mf::Modf => Function::Regular(MODF_FUNCTION),
Mf::Frexp => Function::Regular(FREXP_FUNCTION),
Mf::Ldexp => Function::Regular("ldexp"),
// exponent
Mf::Exp => Function::Regular("exp"),

2
src/back/msl/keywords.rs
View File

@@ -214,4 +214,6 @@ pub const RESERVED: &[&str] = &[
// Naga utilities
"DefaultConstructible",
"clamped_lod_e",
super::writer::FREXP_FUNCTION,
super::writer::MODF_FUNCTION,
];

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

@@ -32,6 +32,9 @@ const RAY_QUERY_FIELD_INTERSECTION: &str = "intersection";
const RAY_QUERY_FIELD_READY: &str = "ready";
const RAY_QUERY_FUN_MAP_INTERSECTION: &str = "_map_intersection_type";
pub(crate) const MODF_FUNCTION: &str = "naga_modf";
pub(crate) const FREXP_FUNCTION: &str = "naga_frexp";
/// Write the Metal name for a Naga numeric type: scalar, vector, or matrix.
///
/// The `sizes` slice determines whether this function writes a
@@ -1678,8 +1681,8 @@ impl<W: Write> Writer<W> {
Mf::Round => "rint",
Mf::Fract => "fract",
Mf::Trunc => "trunc",
Mf::Modf => "modf",
Mf::Frexp => "frexp",
Mf::Modf => MODF_FUNCTION,
Mf::Frexp => FREXP_FUNCTION,
Mf::Ldexp => "ldexp",
// exponent
Mf::Exp => "exp",
@@ -1813,6 +1816,9 @@ impl<W: Write> Writer<W> {
write!(self.out, "((")?;
self.put_expression(arg, context, false)?;
write!(self.out, ") * 57.295779513082322865)")?;
} else if fun == Mf::Modf || fun == Mf::Frexp {
write!(self.out, "{fun_name}")?;
self.put_call_parameters(iter::once(arg), context)?;
} else {
write!(self.out, "{NAMESPACE}::{fun_name}")?;
self.put_call_parameters(
@@ -3236,6 +3242,57 @@ impl<W: Write> Writer<W> {
}
}
}
// Write functions to create special types.
for (type_key, struct_ty) in module.special_types.predeclared_types.iter() {
match type_key {
&crate::PredeclaredType::ModfResult { size, width }
| &crate::PredeclaredType::FrexpResult { size, width } => {
let arg_type_name_owner;
let arg_type_name = if let Some(size) = size {
arg_type_name_owner = format!(
"{NAMESPACE}::{}{}",
if width == 8 { "double" } else { "float" },
size as u8
);
&arg_type_name_owner
} else if width == 8 {
"double"
} else {
"float"
};
let other_type_name_owner;
let (defined_func_name, called_func_name, other_type_name) =
if matches!(type_key, &crate::PredeclaredType::ModfResult { .. }) {
(MODF_FUNCTION, "modf", arg_type_name)
} else {
let other_type_name = if let Some(size) = size {
other_type_name_owner = format!("int{}", size as u8);
&other_type_name_owner
} else {
"int"
};
(FREXP_FUNCTION, "frexp", other_type_name)
};
let struct_name = &self.names[&NameKey::Type(*struct_ty)];
writeln!(self.out)?;
writeln!(
self.out,
"{} {defined_func_name}({arg_type_name} arg) {{
{other_type_name} other;
{arg_type_name} fract = {NAMESPACE}::{called_func_name}(arg, other);
return {}{{ fract, other }};
}}",
struct_name, struct_name
)?;
}
&crate::PredeclaredType::AtomicCompareExchangeWeakResult { .. } => {}
}
}
Ok(())
}

4
src/back/spv/block.rs
View File

@@ -787,8 +787,8 @@ impl<'w> BlockContext<'w> {
Mf::Floor => MathOp::Ext(spirv::GLOp::Floor),
Mf::Fract => MathOp::Ext(spirv::GLOp::Fract),
Mf::Trunc => MathOp::Ext(spirv::GLOp::Trunc),
Mf::Modf => MathOp::Ext(spirv::GLOp::Modf),
Mf::Frexp => MathOp::Ext(spirv::GLOp::Frexp),
Mf::Modf => MathOp::Ext(spirv::GLOp::ModfStruct),
Mf::Frexp => MathOp::Ext(spirv::GLOp::FrexpStruct),
Mf::Ldexp => MathOp::Ext(spirv::GLOp::Ldexp),
// geometry
Mf::Dot => match *self.fun_info[arg].ty.inner_with(&self.ir_module.types) {

22
src/back/wgsl/writer.rs
View File

@@ -97,6 +97,14 @@ impl<W: Write> Writer<W> {
self.ep_results.clear();
}
fn is_builtin_wgsl_struct(&self, module: &Module, handle: Handle<crate::Type>) -> bool {
module
.special_types
.predeclared_types
.values()
.any(|t| *t == handle)
}
pub fn write(&mut self, module: &Module, info: &valid::ModuleInfo) -> BackendResult {
self.reset(module);
@@ -109,13 +117,13 @@ impl<W: Write> Writer<W> {
// Write all structs
for (handle, ty) in module.types.iter() {
if let TypeInner::Struct {
ref members,
span: _,
} = ty.inner
{
self.write_struct(module, handle, members)?;
writeln!(self.out)?;
if let TypeInner::Struct { ref members, .. } = ty.inner {
{
if !self.is_builtin_wgsl_struct(module, handle) {
self.write_struct(module, handle, members)?;
writeln!(self.out)?;
}
}
}
}

248
src/front/type_gen.rs
View File

@@ -5,55 +5,6 @@ Type generators.
use crate::{arena::Handle, span::Span};
impl crate::Module {
pub fn generate_atomic_compare_exchange_result(
&mut self,
kind: crate::ScalarKind,
width: crate::Bytes,
) -> Handle<crate::Type> {
let bool_ty = self.types.insert(
crate::Type {
name: None,
inner: crate::TypeInner::Scalar {
kind: crate::ScalarKind::Bool,
width: crate::BOOL_WIDTH,
},
},
Span::UNDEFINED,
);
let scalar_ty = self.types.insert(
crate::Type {
name: None,
inner: crate::TypeInner::Scalar { kind, width },
},
Span::UNDEFINED,
);
self.types.insert(
crate::Type {
name: Some(format!(
"__atomic_compare_exchange_result<{kind:?},{width}>"
)),
inner: crate::TypeInner::Struct {
members: vec![
crate::StructMember {
name: Some("old_value".to_string()),
ty: scalar_ty,
binding: None,
offset: 0,
},
crate::StructMember {
name: Some("exchanged".to_string()),
ty: bool_ty,
binding: None,
offset: 4,
},
],
span: 8,
},
},
Span::UNDEFINED,
)
}
/// Populate this module's [`SpecialTypes::ray_desc`] type.
///
/// [`SpecialTypes::ray_desc`] is the type of the [`descriptor`] operand of
@@ -311,4 +262,203 @@ impl crate::Module {
self.special_types.ray_intersection = Some(handle);
handle
}
/// Populate this module's [`SpecialTypes::predeclared_types`] type and return the handle.
///
/// [`SpecialTypes::predeclared_types`]: crate::SpecialTypes::predeclared_types
pub fn generate_predeclared_type(
&mut self,
special_type: crate::PredeclaredType,
) -> Handle<crate::Type> {
use std::fmt::Write;
if let Some(value) = self.special_types.predeclared_types.get(&special_type) {
return *value;
}
let ty = match special_type {
crate::PredeclaredType::AtomicCompareExchangeWeakResult { kind, width } => {
let bool_ty = self.types.insert(
crate::Type {
name: None,
inner: crate::TypeInner::Scalar {
kind: crate::ScalarKind::Bool,
width: crate::BOOL_WIDTH,
},
},
Span::UNDEFINED,
);
let scalar_ty = self.types.insert(
crate::Type {
name: None,
inner: crate::TypeInner::Scalar { kind, width },
},
Span::UNDEFINED,
);
crate::Type {
name: Some(format!(
"__atomic_compare_exchange_result<{kind:?},{width}>"
)),
inner: crate::TypeInner::Struct {
members: vec![
crate::StructMember {
name: Some("old_value".to_string()),
ty: scalar_ty,
binding: None,
offset: 0,
},
crate::StructMember {
name: Some("exchanged".to_string()),
ty: bool_ty,
binding: None,
offset: 4,
},
],
span: 8,
},
}
}
crate::PredeclaredType::ModfResult { size, width } => {
let float_ty = self.types.insert(
crate::Type {
name: None,
inner: crate::TypeInner::Scalar {
kind: crate::ScalarKind::Float,
width,
},
},
Span::UNDEFINED,
);
let (member_ty, second_offset) = if let Some(size) = size {
let vec_ty = self.types.insert(
crate::Type {
name: None,
inner: crate::TypeInner::Vector {
size,
kind: crate::ScalarKind::Float,
width,
},
},
Span::UNDEFINED,
);
(vec_ty, size as u32 * width as u32)
} else {
(float_ty, width as u32)
};
let mut type_name = "__modf_result_".to_string();
if let Some(size) = size {
let _ = write!(type_name, "vec{}_", size as u8);
}
let _ = write!(type_name, "f{}", width * 8);
crate::Type {
name: Some(type_name),
inner: crate::TypeInner::Struct {
members: vec![
crate::StructMember {
name: Some("fract".to_string()),
ty: member_ty,
binding: None,
offset: 0,
},
crate::StructMember {
name: Some("whole".to_string()),
ty: member_ty,
binding: None,
offset: second_offset,
},
],
span: second_offset * 2,
},
}
}
crate::PredeclaredType::FrexpResult { size, width } => {
let float_ty = self.types.insert(
crate::Type {
name: None,
inner: crate::TypeInner::Scalar {
kind: crate::ScalarKind::Float,
width,
},
},
Span::UNDEFINED,
);
let int_ty = self.types.insert(
crate::Type {
name: None,
inner: crate::TypeInner::Scalar {
kind: crate::ScalarKind::Sint,
width,
},
},
Span::UNDEFINED,
);
let (fract_member_ty, exp_member_ty, second_offset) = if let Some(size) = size {
let vec_float_ty = self.types.insert(
crate::Type {
name: None,
inner: crate::TypeInner::Vector {
size,
kind: crate::ScalarKind::Float,
width,
},
},
Span::UNDEFINED,
);
let vec_int_ty = self.types.insert(
crate::Type {
name: None,
inner: crate::TypeInner::Vector {
size,
kind: crate::ScalarKind::Sint,
width,
},
},
Span::UNDEFINED,
);
(vec_float_ty, vec_int_ty, size as u32 * width as u32)
} else {
(float_ty, int_ty, width as u32)
};
let mut type_name = "__frexp_result_".to_string();
if let Some(size) = size {
let _ = write!(type_name, "vec{}_", size as u8);
}
let _ = write!(type_name, "f{}", width * 8);
crate::Type {
name: Some(type_name),
inner: crate::TypeInner::Struct {
members: vec![
crate::StructMember {
name: Some("fract".to_string()),
ty: fract_member_ty,
binding: None,
offset: 0,
},
crate::StructMember {
name: Some("exp".to_string()),
ty: exp_member_ty,
binding: None,
offset: second_offset,
},
],
span: second_offset * 2,
},
}
}
};
let handle = self.types.insert(ty, Span::UNDEFINED);
self.special_types
.predeclared_types
.insert(special_type, handle);
handle
}
}

29
src/front/wgsl/lower/mod.rs
View File

@@ -1745,6 +1745,25 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
args.finish()?;
if fun == crate::MathFunction::Modf || fun == crate::MathFunction::Frexp {
ctx.grow_types(arg)?;
if let Some((size, width)) = match *ctx.resolved_inner(arg) {
crate::TypeInner::Scalar { width, .. } => Some((None, width)),
crate::TypeInner::Vector { size, width, .. } => {
Some((Some(size), width))
}
_ => None,
} {
ctx.module.generate_predeclared_type(
if fun == crate::MathFunction::Modf {
crate::PredeclaredType::ModfResult { size, width }
} else {
crate::PredeclaredType::FrexpResult { size, width }
},
);
}
}
crate::Expression::Math {
fun,
arg,
@@ -1880,10 +1899,12 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
let expression = match *ctx.resolved_inner(value) {
crate::TypeInner::Scalar { kind, width } => {
crate::Expression::AtomicResult {
//TODO: cache this to avoid generating duplicate types
ty: ctx
.module
.generate_atomic_compare_exchange_result(kind, width),
ty: ctx.module.generate_predeclared_type(
crate::PredeclaredType::AtomicCompareExchangeWeakResult {
kind,
width,
},
),
comparison: true,
}
}

5
src/front/wgsl/tests.rs
View File

@@ -456,10 +456,11 @@ fn binary_expression_mixed_scalar_and_vector_operands() {
#[test]
fn parse_pointers() {
parse_str(
"fn foo() {
"fn foo(a: ptr<private, f32>) -> f32 { return *a; }
fn bar() {
var x: f32 = 1.0;
let px = &x;
let py = frexp(0.5, px);
let py = foo(px);
}",
)
.unwrap();

31
src/lib.rs
View File

@@ -1945,6 +1945,31 @@ pub struct EntryPoint {
pub function: Function,
}
/// Return types predeclared for the frexp, modf, and atomicCompareExchangeWeak built-in functions.
///
/// These cannot be spelled in WGSL source.
///
/// Stored in [`SpecialTypes::predeclared_types`] and created by [`Module::generate_predeclared_type`].
#[derive(Debug, PartialEq, Eq, Hash)]
#[cfg_attr(feature = "clone", derive(Clone))]
#[cfg_attr(feature = "serialize", derive(Serialize))]
#[cfg_attr(feature = "deserialize", derive(Deserialize))]
#[cfg_attr(feature = "arbitrary", derive(Arbitrary))]
pub enum PredeclaredType {
AtomicCompareExchangeWeakResult {
kind: ScalarKind,
width: Bytes,
},
ModfResult {
size: Option<VectorSize>,
width: Bytes,
},
FrexpResult {
size: Option<VectorSize>,
width: Bytes,
},
}
/// Set of special types that can be optionally generated by the frontends.
#[derive(Debug, Default)]
#[cfg_attr(feature = "clone", derive(Clone))]
@@ -1963,6 +1988,12 @@ pub struct SpecialTypes {
/// Call [`Module::generate_ray_intersection_type`] to populate
/// this if needed and return the handle.
pub ray_intersection: Option<Handle<Type>>,
/// Types for predeclared wgsl types instantiated on demand.
///
/// Call [`Module::generate_predeclared_type`] to populate this if
/// needed and return the handle.
pub predeclared_types: indexmap::IndexMap<PredeclaredType, Handle<Type>>,
}
/// Shader module.

4
src/proc/mod.rs
View File

@@ -375,8 +375,8 @@ impl super::MathFunction {
Self::Round => 1,
Self::Fract => 1,
Self::Trunc => 1,
Self::Modf => 2,
Self::Frexp => 2,
Self::Modf => 1,
Self::Frexp => 1,
Self::Ldexp => 2,
// exponent
Self::Exp => 1,

27
src/proc/typifier.rs
View File

@@ -706,8 +706,6 @@ impl<'a> ResolveContext<'a> {
Mf::Round |
Mf::Fract |
Mf::Trunc |
Mf::Modf |
Mf::Frexp |
Mf::Ldexp |
// exponent
Mf::Exp |
@@ -715,6 +713,31 @@ impl<'a> ResolveContext<'a> {
Mf::Log |
Mf::Log2 |
Mf::Pow => res_arg.clone(),
Mf::Modf | Mf::Frexp => {
let (size, width) = match res_arg.inner_with(types) {
&Ti::Scalar {
kind: crate::ScalarKind::Float,
width,
} => (None, width),
&Ti::Vector {
kind: crate::ScalarKind::Float,
size,
width,
} => (Some(size), width),
ref other =>
return Err(ResolveError::IncompatibleOperands(format!("{fun:?}({other:?}, _)")))
};
let result = self
.special_types
.predeclared_types
.get(&if fun == Mf::Modf {
crate::PredeclaredType::ModfResult { size, width }
} else {
crate::PredeclaredType::FrexpResult { size, width }
})
.ok_or(ResolveError::MissingSpecialType)?;
TypeResolution::Handle(*result)
},
// geometry
Mf::Dot => match *res_arg.inner_with(types) {
Ti::Vector {

40
src/valid/expression.rs
View File

@@ -1015,38 +1015,20 @@ impl super::Validator {
}
}
Mf::Modf | Mf::Frexp => {
let arg1_ty = match (arg1_ty, arg2_ty, arg3_ty) {
(Some(ty1), None, None) => ty1,
_ => return Err(ExpressionError::WrongArgumentCount(fun)),
};
let (size0, width0) = match *arg_ty {
if arg1_ty.is_some() | arg2_ty.is_some() | arg3_ty.is_some() {
return Err(ExpressionError::WrongArgumentCount(fun));
}
if !matches!(
*arg_ty,
Ti::Scalar {
kind: Sk::Float,
width,
} => (None, width),
Ti::Vector {
..
} | Ti::Vector {
kind: Sk::Float,
size,
width,
} => (Some(size), width),
_ => return Err(ExpressionError::InvalidArgumentType(fun, 0, arg)),
};
let good = match *arg1_ty {
Ti::Pointer { base, space: _ } => module.types[base].inner == *arg_ty,
Ti::ValuePointer {
size,
kind: Sk::Float,
width,
space: _,
} => size == size0 && width == width0,
_ => false,
};
if !good {
return Err(ExpressionError::InvalidArgumentType(
fun,
1,
arg1.unwrap(),
));
..
},
) {
return Err(ExpressionError::InvalidArgumentType(fun, 1, arg));
}
}
Mf::Ldexp => {

10
tests/in/math-functions.wgsl
View File

@@ -31,4 +31,14 @@ fn main() {
let clz_d = countLeadingZeros(vec2(1u));
let lde_a = ldexp(1.0, 2);
let lde_b = ldexp(vec2(1.0, 2.0), vec2(3, 4));
let modf_a = modf(1.5);
let modf_b = modf(1.5).fract;
let modf_c = modf(1.5).whole;
let modf_d = modf(vec2(1.5, 1.5));
let modf_e = modf(vec4(1.5, 1.5, 1.5, 1.5)).whole.x;
let modf_f: f32 = modf(vec2(1.5, 1.5)).fract.y;
let frexp_a = frexp(1.5);
let frexp_b = frexp(1.5).fract;
let frexp_c: i32 = frexp(1.5).exp;
let frexp_d: i32 = frexp(vec4(1.5, 1.5, 1.5, 1.5)).exp.x;
}

60
tests/out/glsl/math-functions.main.Fragment.glsl
View File

@@ -3,6 +3,56 @@
precision highp float;
precision highp int;
struct __modf_result_f32_ {
float fract_;
float whole;
};
struct __modf_result_vec2_f32_ {
vec2 fract_;
vec2 whole;
};
struct __modf_result_vec4_f32_ {
vec4 fract_;
vec4 whole;
};
struct __frexp_result_f32_ {
float fract_;
int exp_;
};
struct __frexp_result_vec4_f32_ {
vec4 fract_;
ivec4 exp_;
};
__modf_result_f32_ naga_modf(float arg) {
float other;
float fract = modf(arg, other);
return __modf_result_f32_(fract, other);
}
__modf_result_vec2_f32_ naga_modf(vec2 arg) {
vec2 other;
vec2 fract = modf(arg, other);
return __modf_result_vec2_f32_(fract, other);
}
__modf_result_vec4_f32_ naga_modf(vec4 arg) {
vec4 other;
vec4 fract = modf(arg, other);
return __modf_result_vec4_f32_(fract, other);
}
__frexp_result_f32_ naga_frexp(float arg) {
int other;
float fract = frexp(arg, other);
return __frexp_result_f32_(fract, other);
}
__frexp_result_vec4_f32_ naga_frexp(vec4 arg) {
ivec4 other;
vec4 fract = frexp(arg, other);
return __frexp_result_vec4_f32_(fract, other);
}
void main() {
vec4 v = vec4(0.0);
@@ -36,5 +86,15 @@ void main() {
uvec2 clz_d = uvec2(ivec2(31) - findMSB(uvec2(1u)));
float lde_a = ldexp(1.0, 2);
vec2 lde_b = ldexp(vec2(1.0, 2.0), ivec2(3, 4));
__modf_result_f32_ modf_a = naga_modf(1.5);
float modf_b = naga_modf(1.5).fract_;
float modf_c = naga_modf(1.5).whole;
__modf_result_vec2_f32_ modf_d = naga_modf(vec2(1.5, 1.5));
float modf_e = naga_modf(vec4(1.5, 1.5, 1.5, 1.5)).whole.x;
float modf_f = naga_modf(vec2(1.5, 1.5)).fract_.y;
__frexp_result_f32_ frexp_a = naga_frexp(1.5);
float frexp_b = naga_frexp(1.5).fract_;
int frexp_c = naga_frexp(1.5).exp_;
int frexp_d = naga_frexp(vec4(1.5, 1.5, 1.5, 1.5)).exp_.x;
}

75
tests/out/hlsl/math-functions.hlsl
View File

@@ -1,3 +1,68 @@
struct __modf_result_f32_ {
float fract;
float whole;
};
struct __modf_result_vec2_f32_ {
float2 fract;
float2 whole;
};
struct __modf_result_vec4_f32_ {
float4 fract;
float4 whole;
};
struct __frexp_result_f32_ {
float fract;
int exp_;
};
struct __frexp_result_vec4_f32_ {
float4 fract;
int4 exp_;
};
__modf_result_f32_ naga_modf(float arg) {
float other;
__modf_result_f32_ result;
result.fract = modf(arg, other);
result.whole = other;
return result;
}
__modf_result_vec2_f32_ naga_modf(float2 arg) {
float2 other;
__modf_result_vec2_f32_ result;
result.fract = modf(arg, other);
result.whole = other;
return result;
}
__modf_result_vec4_f32_ naga_modf(float4 arg) {
float4 other;
__modf_result_vec4_f32_ result;
result.fract = modf(arg, other);
result.whole = other;
return result;
}
__frexp_result_f32_ naga_frexp(float arg) {
float other;
__frexp_result_f32_ result;
result.fract = sign(arg) * frexp(arg, other);
result.exp_ = other;
return result;
}
__frexp_result_vec4_f32_ naga_frexp(float4 arg) {
float4 other;
__frexp_result_vec4_f32_ result;
result.fract = sign(arg) * frexp(arg, other);
result.exp_ = other;
return result;
}
void main()
{
float4 v = (0.0).xxxx;
@@ -31,4 +96,14 @@ void main()
uint2 clz_d = ((31u).xx - firstbithigh((1u).xx));
float lde_a = ldexp(1.0, 2);
float2 lde_b = ldexp(float2(1.0, 2.0), int2(3, 4));
__modf_result_f32_ modf_a = naga_modf(1.5);
float modf_b = naga_modf(1.5).fract;
float modf_c = naga_modf(1.5).whole;
__modf_result_vec2_f32_ modf_d = naga_modf(float2(1.5, 1.5));
float modf_e = naga_modf(float4(1.5, 1.5, 1.5, 1.5)).whole.x;
float modf_f = naga_modf(float2(1.5, 1.5)).fract.y;
__frexp_result_f32_ frexp_a = naga_frexp(1.5);
float frexp_b = naga_frexp(1.5).fract;
int frexp_c = naga_frexp(1.5).exp_;
int frexp_d = naga_frexp(float4(1.5, 1.5, 1.5, 1.5)).exp_.x;
}

1
tests/out/ir/access.ron
View File

@@ -334,6 +334,7 @@
special_types: (
ray_desc: None,
ray_intersection: None,
predeclared_types: {},
),
constants: [],
global_variables: [

1
tests/out/ir/collatz.ron
View File

@@ -41,6 +41,7 @@
special_types: (
ray_desc: None,
ray_intersection: None,
predeclared_types: {},
),
constants: [],
global_variables: [

1
tests/out/ir/shadow.ron
View File

@@ -266,6 +266,7 @@
special_types: (
ray_desc: None,
ray_intersection: None,
predeclared_types: {},
),
constants: [
(

60
tests/out/msl/math-functions.msl
View File

@@ -4,6 +4,56 @@
using metal::uint;
struct __modf_result_f32_ {
float fract;
float whole;
};
struct __modf_result_vec2_f32_ {
metal::float2 fract;
metal::float2 whole;
};
struct __modf_result_vec4_f32_ {
metal::float4 fract;
metal::float4 whole;
};
struct __frexp_result_f32_ {
float fract;
int exp;
};
struct __frexp_result_vec4_f32_ {
metal::float4 fract;
metal::int4 exp;
};
__modf_result_f32_ naga_modf(float arg) {
float other;
float fract = metal::modf(arg, other);
return __modf_result_f32_{ fract, other };
}
__modf_result_vec2_f32_ naga_modf(metal::float2 arg) {
metal::float2 other;
metal::float2 fract = metal::modf(arg, other);
return __modf_result_vec2_f32_{ fract, other };
}
__modf_result_vec4_f32_ naga_modf(metal::float4 arg) {
metal::float4 other;
metal::float4 fract = metal::modf(arg, other);
return __modf_result_vec4_f32_{ fract, other };
}
__frexp_result_f32_ naga_frexp(float arg) {
int other;
float fract = metal::frexp(arg, other);
return __frexp_result_f32_{ fract, other };
}
__frexp_result_vec4_f32_ naga_frexp(metal::float4 arg) {
int4 other;
metal::float4 fract = metal::frexp(arg, other);
return __frexp_result_vec4_f32_{ fract, other };
}
fragment void main_(
) {
@@ -40,4 +90,14 @@ fragment void main_(
metal::uint2 clz_d = metal::clz(metal::uint2(1u));
float lde_a = metal::ldexp(1.0, 2);
metal::float2 lde_b = metal::ldexp(metal::float2(1.0, 2.0), metal::int2(3, 4));
__modf_result_f32_ modf_a = naga_modf(1.5);
float modf_b = naga_modf(1.5).fract;
float modf_c = naga_modf(1.5).whole;
__modf_result_vec2_f32_ modf_d = naga_modf(metal::float2(1.5, 1.5));
float modf_e = naga_modf(metal::float4(1.5, 1.5, 1.5, 1.5)).whole.x;
float modf_f = naga_modf(metal::float2(1.5, 1.5)).fract.y;
__frexp_result_f32_ frexp_a = naga_frexp(1.5);
float frexp_b = naga_frexp(1.5).fract;
int frexp_c = naga_frexp(1.5).exp;
int frexp_d = naga_frexp(metal::float4(1.5, 1.5, 1.5, 1.5)).exp.x;
}

225
tests/out/spv/math-functions.spvasm
View File

@@ -1,106 +1,147 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 96
; Bound: 127
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %9 "main"
OpExecutionMode %9 OriginUpperLeft
OpEntryPoint Fragment %15 "main"
OpExecutionMode %15 OriginUpperLeft
OpMemberDecorate %8 0 Offset 0
OpMemberDecorate %8 1 Offset 4
OpMemberDecorate %9 0 Offset 0
OpMemberDecorate %9 1 Offset 8
OpMemberDecorate %10 0 Offset 0
OpMemberDecorate %10 1 Offset 16
OpMemberDecorate %11 0 Offset 0
OpMemberDecorate %11 1 Offset 4
OpMemberDecorate %13 0 Offset 0
OpMemberDecorate %13 1 Offset 16
%2 = OpTypeVoid
%4 = OpTypeFloat 32
%3 = OpTypeVector %4 4
%6 = OpTypeInt 32 1
%5 = OpTypeVector %6 2
%7 = OpTypeVector %4 2
%10 = OpTypeFunction %2
%11 = OpConstant %4 1.0
%12 = OpConstant %4 0.0
%13 = OpConstantNull %5
%14 = OpTypeInt 32 0
%15 = OpConstant %14 0
%16 = OpConstant %6 -1
%17 = OpConstant %14 1
%18 = OpConstant %6 0
%19 = OpConstant %14 4294967295
%20 = OpConstant %6 1
%21 = OpConstant %6 2
%22 = OpConstant %4 2.0
%23 = OpConstant %6 3
%24 = OpConstant %6 4
%32 = OpConstantComposite %3 %12 %12 %12 %12
%33 = OpConstantComposite %3 %11 %11 %11 %11
%36 = OpConstantNull %6
%49 = OpTypeVector %14 2
%59 = OpConstant %14 32
%69 = OpConstantComposite %49 %59 %59
%81 = OpConstant %6 31
%87 = OpConstantComposite %5 %81 %81
%9 = OpFunction %2 None %10
%8 = OpLabel
OpBranch %25
%25 = OpLabel
%26 = OpCompositeConstruct %3 %12 %12 %12 %12
%27 = OpExtInst %4 %1 Degrees %11
%28 = OpExtInst %4 %1 Radians %11
%29 = OpExtInst %3 %1 Degrees %26
%30 = OpExtInst %3 %1 Radians %26
%31 = OpExtInst %3 %1 FClamp %26 %32 %33
%34 = OpExtInst %3 %1 Refract %26 %26 %11
%37 = OpCompositeExtract %6 %13 0
%38 = OpCompositeExtract %6 %13 0
%39 = OpIMul %6 %37 %38
%40 = OpIAdd %6 %36 %39
%41 = OpCompositeExtract %6 %13 1
%42 = OpCompositeExtract %6 %13 1
%43 = OpIMul %6 %41 %42
%35 = OpIAdd %6 %40 %43
%44 = OpCopyObject %14 %15
%45 = OpExtInst %14 %1 FindUMsb %44
%46 = OpExtInst %6 %1 FindSMsb %16
%47 = OpCompositeConstruct %5 %16 %16
%48 = OpExtInst %5 %1 FindSMsb %47
%50 = OpCompositeConstruct %49 %17 %17
%51 = OpExtInst %49 %1 FindUMsb %50
%52 = OpExtInst %6 %1 FindILsb %16
%53 = OpExtInst %14 %1 FindILsb %17
%54 = OpCompositeConstruct %5 %16 %16
%55 = OpExtInst %5 %1 FindILsb %54
%56 = OpCompositeConstruct %49 %17 %17
%57 = OpExtInst %49 %1 FindILsb %56
%60 = OpExtInst %14 %1 FindILsb %15
%58 = OpExtInst %14 %1 UMin %59 %60
%62 = OpExtInst %6 %1 FindILsb %18
%61 = OpExtInst %6 %1 UMin %59 %62
%64 = OpExtInst %14 %1 FindILsb %19
%63 = OpExtInst %14 %1 UMin %59 %64
%66 = OpExtInst %6 %1 FindILsb %16
%65 = OpExtInst %6 %1 UMin %59 %66
%67 = OpCompositeConstruct %49 %15 %15
%70 = OpExtInst %49 %1 FindILsb %67
%68 = OpExtInst %49 %1 UMin %69 %70
%71 = OpCompositeConstruct %5 %18 %18
%73 = OpExtInst %5 %1 FindILsb %71
%72 = OpExtInst %5 %1 UMin %69 %73
%74 = OpCompositeConstruct %49 %17 %17
%76 = OpExtInst %49 %1 FindILsb %74
%75 = OpExtInst %49 %1 UMin %69 %76
%77 = OpCompositeConstruct %5 %20 %20
%79 = OpExtInst %5 %1 FindILsb %77
%78 = OpExtInst %5 %1 UMin %69 %79
%82 = OpExtInst %6 %1 FindUMsb %16
%80 = OpISub %6 %81 %82
%84 = OpExtInst %6 %1 FindUMsb %17
%83 = OpISub %14 %81 %84
%85 = OpCompositeConstruct %5 %16 %16
%88 = OpExtInst %5 %1 FindUMsb %85
%86 = OpISub %5 %87 %88
%89 = OpCompositeConstruct %49 %17 %17
%91 = OpExtInst %5 %1 FindUMsb %89
%90 = OpISub %49 %87 %91
%92 = OpExtInst %4 %1 Ldexp %11 %21
%93 = OpCompositeConstruct %7 %11 %22
%94 = OpCompositeConstruct %5 %23 %24
%95 = OpExtInst %7 %1 Ldexp %93 %94
%8 = OpTypeStruct %4 %4
%9 = OpTypeStruct %7 %7
%10 = OpTypeStruct %3 %3
%11 = OpTypeStruct %4 %6
%12 = OpTypeVector %6 4
%13 = OpTypeStruct %3 %12
%16 = OpTypeFunction %2
%17 = OpConstant %4 1.0
%18 = OpConstant %4 0.0
%19 = OpConstantNull %5
%20 = OpTypeInt 32 0
%21 = OpConstant %20 0
%22 = OpConstant %6 -1
%23 = OpConstant %20 1
%24 = OpConstant %6 0
%25 = OpConstant %20 4294967295
%26 = OpConstant %6 1
%27 = OpConstant %6 2
%28 = OpConstant %4 2.0
%29 = OpConstant %6 3
%30 = OpConstant %6 4
%31 = OpConstant %4 1.5
%39 = OpConstantComposite %3 %18 %18 %18 %18
%40 = OpConstantComposite %3 %17 %17 %17 %17
%43 = OpConstantNull %6
%56 = OpTypeVector %20 2
%66 = OpConstant %20 32
%76 = OpConstantComposite %56 %66 %66
%88 = OpConstant %6 31
%94 = OpConstantComposite %5 %88 %88
%15 = OpFunction %2 None %16
%14 = OpLabel
OpBranch %32
%32 = OpLabel
%33 = OpCompositeConstruct %3 %18 %18 %18 %18
%34 = OpExtInst %4 %1 Degrees %17
%35 = OpExtInst %4 %1 Radians %17
%36 = OpExtInst %3 %1 Degrees %33
%37 = OpExtInst %3 %1 Radians %33
%38 = OpExtInst %3 %1 FClamp %33 %39 %40
%41 = OpExtInst %3 %1 Refract %33 %33 %17
%44 = OpCompositeExtract %6 %19 0
%45 = OpCompositeExtract %6 %19 0
%46 = OpIMul %6 %44 %45
%47 = OpIAdd %6 %43 %46
%48 = OpCompositeExtract %6 %19 1
%49 = OpCompositeExtract %6 %19 1
%50 = OpIMul %6 %48 %49
%42 = OpIAdd %6 %47 %50
%51 = OpCopyObject %20 %21
%52 = OpExtInst %20 %1 FindUMsb %51
%53 = OpExtInst %6 %1 FindSMsb %22
%54 = OpCompositeConstruct %5 %22 %22
%55 = OpExtInst %5 %1 FindSMsb %54
%57 = OpCompositeConstruct %56 %23 %23
%58 = OpExtInst %56 %1 FindUMsb %57
%59 = OpExtInst %6 %1 FindILsb %22
%60 = OpExtInst %20 %1 FindILsb %23
%61 = OpCompositeConstruct %5 %22 %22
%62 = OpExtInst %5 %1 FindILsb %61
%63 = OpCompositeConstruct %56 %23 %23
%64 = OpExtInst %56 %1 FindILsb %63
%67 = OpExtInst %20 %1 FindILsb %21
%65 = OpExtInst %20 %1 UMin %66 %67
%69 = OpExtInst %6 %1 FindILsb %24
%68 = OpExtInst %6 %1 UMin %66 %69
%71 = OpExtInst %20 %1 FindILsb %25
%70 = OpExtInst %20 %1 UMin %66 %71
%73 = OpExtInst %6 %1 FindILsb %22
%72 = OpExtInst %6 %1 UMin %66 %73
%74 = OpCompositeConstruct %56 %21 %21
%77 = OpExtInst %56 %1 FindILsb %74
%75 = OpExtInst %56 %1 UMin %76 %77
%78 = OpCompositeConstruct %5 %24 %24
%80 = OpExtInst %5 %1 FindILsb %78
%79 = OpExtInst %5 %1 UMin %76 %80
%81 = OpCompositeConstruct %56 %23 %23
%83 = OpExtInst %56 %1 FindILsb %81
%82 = OpExtInst %56 %1 UMin %76 %83
%84 = OpCompositeConstruct %5 %26 %26
%86 = OpExtInst %5 %1 FindILsb %84
%85 = OpExtInst %5 %1 UMin %76 %86
%89 = OpExtInst %6 %1 FindUMsb %22
%87 = OpISub %6 %88 %89
%91 = OpExtInst %6 %1 FindUMsb %23
%90 = OpISub %20 %88 %91
%92 = OpCompositeConstruct %5 %22 %22
%95 = OpExtInst %5 %1 FindUMsb %92
%93 = OpISub %5 %94 %95
%96 = OpCompositeConstruct %56 %23 %23
%98 = OpExtInst %5 %1 FindUMsb %96
%97 = OpISub %56 %94 %98
%99 = OpExtInst %4 %1 Ldexp %17 %27
%100 = OpCompositeConstruct %7 %17 %28
%101 = OpCompositeConstruct %5 %29 %30
%102 = OpExtInst %7 %1 Ldexp %100 %101
%103 = OpExtInst %8 %1 ModfStruct %31
%104 = OpExtInst %8 %1 ModfStruct %31
%105 = OpCompositeExtract %4 %104 0
%106 = OpExtInst %8 %1 ModfStruct %31
%107 = OpCompositeExtract %4 %106 1
%108 = OpCompositeConstruct %7 %31 %31
%109 = OpExtInst %9 %1 ModfStruct %108
%110 = OpCompositeConstruct %3 %31 %31 %31 %31
%111 = OpExtInst %10 %1 ModfStruct %110
%112 = OpCompositeExtract %3 %111 1
%113 = OpCompositeExtract %4 %112 0
%114 = OpCompositeConstruct %7 %31 %31
%115 = OpExtInst %9 %1 ModfStruct %114
%116 = OpCompositeExtract %7 %115 0
%117 = OpCompositeExtract %4 %116 1
%118 = OpExtInst %11 %1 FrexpStruct %31
%119 = OpExtInst %11 %1 FrexpStruct %31
%120 = OpCompositeExtract %4 %119 0
%121 = OpExtInst %11 %1 FrexpStruct %31
%122 = OpCompositeExtract %6 %121 1
%123 = OpCompositeConstruct %3 %31 %31 %31 %31
%124 = OpExtInst %13 %1 FrexpStruct %123
%125 = OpCompositeExtract %12 %124 1
%126 = OpCompositeExtract %6 %125 0
OpReturn
OpFunctionEnd

10
tests/out/wgsl/atomicCompareExchange.wgsl
View File

@@ -1,13 +1,3 @@
struct gen___atomic_compare_exchange_resultSint4_ {
old_value: i32,
exchanged: bool,
}
struct gen___atomic_compare_exchange_resultUint4_ {
old_value: u32,
exchanged: bool,
}
const SIZE: u32 = 128u;
@group(0) @binding(0)

10
tests/out/wgsl/math-functions.wgsl
View File

@@ -30,4 +30,14 @@ fn main() {
let clz_d = countLeadingZeros(vec2<u32>(1u));
let lde_a = ldexp(1.0, 2);
let lde_b = ldexp(vec2<f32>(1.0, 2.0), vec2<i32>(3, 4));
let modf_a = modf(1.5);
let modf_b = modf(1.5).fract;
let modf_c = modf(1.5).whole;
let modf_d = modf(vec2<f32>(1.5, 1.5));
let modf_e = modf(vec4<f32>(1.5, 1.5, 1.5, 1.5)).whole.x;
let modf_f = modf(vec2<f32>(1.5, 1.5)).fract.y;
let frexp_a = frexp(1.5);
let frexp_b = frexp(1.5).fract;
let frexp_c = frexp(1.5).exp;
let frexp_d = frexp(vec4<f32>(1.5, 1.5, 1.5, 1.5)).exp.x;
}