diff --git a/src/back/glsl/keywords.rs b/src/back/glsl/keywords.rs index 9679020a0a..afadd6e7f1 100644 --- a/src/back/glsl/keywords.rs +++ b/src/back/glsl/keywords.rs @@ -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, ]; diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index ef5dd0143f..9449cf28f3 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -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; @@ -625,6 +628,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 @@ -2985,8 +3035,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", diff --git a/src/back/hlsl/help.rs b/src/back/hlsl/help.rs index 7ad4631315..2d725514b2 100644 --- a/src/back/hlsl/help.rs +++ b/src/back/hlsl/help.rs @@ -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, diff --git a/src/back/hlsl/keywords.rs b/src/back/hlsl/keywords.rs index 81b797bbf5..059e533ff7 100644 --- a/src/back/hlsl/keywords.rs +++ b/src/back/hlsl/keywords.rs @@ -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 diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 4f19126388..c4cc0bcdce 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -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, @@ -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 @@ -2665,8 +2670,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"), diff --git a/src/back/msl/keywords.rs b/src/back/msl/keywords.rs index a3a9c52dcc..f34b618db8 100644 --- a/src/back/msl/keywords.rs +++ b/src/back/msl/keywords.rs @@ -214,4 +214,6 @@ pub const RESERVED: &[&str] = &[ // Naga utilities "DefaultConstructible", "clamped_lod_e", + super::writer::FREXP_FUNCTION, + super::writer::MODF_FUNCTION, ]; diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index a8a103e6d0..24c54ff8ab 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -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 Writer { 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 Writer { 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 Writer { } } } + + // 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(()) } diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index 11841bc522..c5246ad190 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -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) { diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index 44dd2a97c0..9da6735571 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -96,6 +96,13 @@ impl Writer { self.ep_results.clear(); } + fn is_builtin_wgsl_struct(&self, ty_struct: &crate::Type) -> bool { + ty_struct + .name + .as_ref() + .map_or(false, |name| name.starts_with("__")) + } + pub fn write(&mut self, module: &Module, info: &valid::ModuleInfo) -> BackendResult { self.reset(module); @@ -108,13 +115,13 @@ impl Writer { // 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(ty) { + self.write_struct(module, handle, members)?; + writeln!(self.out)?; + } + } } } diff --git a/src/front/type_gen.rs b/src/front/type_gen.rs index 1ee454c448..0850b7d17b 100644 --- a/src/front/type_gen.rs +++ b/src/front/type_gen.rs @@ -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 { - 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 { + 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, 4) + }; + + 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: 4, + }, + }, + Span::UNDEFINED, + ); + (vec_float_ty, vec_int_ty, size as u32 * width as u32) + } else { + (float_ty, int_ty, 4) + }; + + 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 + } } diff --git a/src/front/wgsl/lower/mod.rs b/src/front/wgsl/lower/mod.rs index 4c5acf0ad9..0ba08d08ff 100644 --- a/src/front/wgsl/lower/mod.rs +++ b/src/front/wgsl/lower/mod.rs @@ -1726,6 +1726,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, @@ -1861,10 +1880,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, } } diff --git a/src/front/wgsl/tests.rs b/src/front/wgsl/tests.rs index 09a4356b07..828f1e3c71 100644 --- a/src/front/wgsl/tests.rs +++ b/src/front/wgsl/tests.rs @@ -456,10 +456,11 @@ fn binary_expression_mixed_scalar_and_vector_operands() { #[test] fn parse_pointers() { parse_str( - "fn foo() { + "fn foo(a: ptr) -> f32 { return *a; } + fn bar() { var x: f32 = 1.0; let px = &x; - let py = frexp(0.5, px); + let py = foo(px); }", ) .unwrap(); diff --git a/src/lib.rs b/src/lib.rs index a1f7ef1654..c90eab6b85 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -1943,6 +1943,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, + width: Bytes, + }, + FrexpResult { + size: Option, + width: Bytes, + }, +} + /// Set of special types that can be optionally generated by the frontends. #[derive(Debug, Default)] #[cfg_attr(feature = "clone", derive(Clone))] @@ -1961,6 +1986,12 @@ pub struct SpecialTypes { /// Call [`Module::generate_ray_intersection_type`] to populate /// this if needed and return the handle. pub ray_intersection: Option>, + + /// 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>, } /// Shader module. diff --git a/src/proc/mod.rs b/src/proc/mod.rs index 582c887034..bfb2b0d7ac 100644 --- a/src/proc/mod.rs +++ b/src/proc/mod.rs @@ -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, diff --git a/src/proc/typifier.rs b/src/proc/typifier.rs index 6b79e0ead2..a6130ad796 100644 --- a/src/proc/typifier.rs +++ b/src/proc/typifier.rs @@ -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 { diff --git a/src/valid/expression.rs b/src/valid/expression.rs index 1703d213e1..7025f1ad92 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -1015,38 +1015,20 @@ impl super::Validator { } } Mf::Modf | Mf::Frexp | Mf::Ldexp => { - 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 { - 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, + .. + } | Ti::Vector { 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::Dot => { diff --git a/tests/in/math-functions.wgsl b/tests/in/math-functions.wgsl index efeb988f5a..d3c6f4fca9 100644 --- a/tests/in/math-functions.wgsl +++ b/tests/in/math-functions.wgsl @@ -29,4 +29,14 @@ fn main() { let clz_b = countLeadingZeros(1u); let clz_c = countLeadingZeros(vec2(-1)); let clz_d = countLeadingZeros(vec2(1u)); + 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; } diff --git a/tests/out/glsl/math-functions.main.Fragment.glsl b/tests/out/glsl/math-functions.main.Fragment.glsl index 70ad0b199a..128bbbc89a 100644 --- a/tests/out/glsl/math-functions.main.Fragment.glsl +++ b/tests/out/glsl/math-functions.main.Fragment.glsl @@ -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); @@ -34,5 +84,15 @@ void main() { ivec2 _e58 = ivec2(-1); ivec2 clz_c = mix(ivec2(31) - findMSB(_e58), ivec2(0), lessThan(_e58, ivec2(0))); uvec2 clz_d = uvec2(ivec2(31) - findMSB(uvec2(1u))); + __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; } diff --git a/tests/out/hlsl/math-functions.hlsl b/tests/out/hlsl/math-functions.hlsl index a1be810532..f0e46addf4 100644 --- a/tests/out/hlsl/math-functions.hlsl +++ b/tests/out/hlsl/math-functions.hlsl @@ -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; @@ -29,4 +94,14 @@ void main() int2 _expr58 = (-1).xx; int2 clz_c = (_expr58 < (0).xx ? (0).xx : (31).xx - asint(firstbithigh(_expr58))); uint2 clz_d = ((31u).xx - firstbithigh((1u).xx)); + __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; } diff --git a/tests/out/ir/access.ron b/tests/out/ir/access.ron index 7447249127..70a86f44dc 100644 --- a/tests/out/ir/access.ron +++ b/tests/out/ir/access.ron @@ -334,6 +334,7 @@ special_types: ( ray_desc: None, ray_intersection: None, + special_types: {}, ), constants: [], global_variables: [ diff --git a/tests/out/ir/collatz.ron b/tests/out/ir/collatz.ron index 6451fc715f..6eb879588f 100644 --- a/tests/out/ir/collatz.ron +++ b/tests/out/ir/collatz.ron @@ -41,6 +41,7 @@ special_types: ( ray_desc: None, ray_intersection: None, + special_types: {}, ), constants: [], global_variables: [ diff --git a/tests/out/ir/shadow.ron b/tests/out/ir/shadow.ron index 246657255b..55b850b07c 100644 --- a/tests/out/ir/shadow.ron +++ b/tests/out/ir/shadow.ron @@ -266,6 +266,7 @@ special_types: ( ray_desc: None, ray_intersection: None, + special_types: {}, ), constants: [ ( diff --git a/tests/out/msl/math-functions.msl b/tests/out/msl/math-functions.msl index 120d526d44..08de0160a5 100644 --- a/tests/out/msl/math-functions.msl +++ b/tests/out/msl/math-functions.msl @@ -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_( ) { @@ -38,4 +88,14 @@ fragment void main_( uint clz_b = metal::clz(1u); metal::int2 clz_c = metal::clz(metal::int2(-1)); metal::uint2 clz_d = metal::clz(metal::uint2(1u)); + __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; } diff --git a/tests/out/spv/math-functions.spvasm b/tests/out/spv/math-functions.spvasm index 60bdd7ae0d..61d2b52b52 100644 --- a/tests/out/spv/math-functions.spvasm +++ b/tests/out/spv/math-functions.spvasm @@ -1,97 +1,139 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 87 +; Bound: 119 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Fragment %8 "main" -OpExecutionMode %8 OriginUpperLeft +OpEntryPoint Fragment %15 "main" +OpExecutionMode %15 OriginUpperLeft +OpMemberDecorate %7 0 Offset 0 +OpMemberDecorate %7 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 -%9 = OpTypeFunction %2 -%10 = OpConstant %4 1.0 -%11 = OpConstant %4 0.0 -%12 = OpConstantNull %5 -%13 = OpTypeInt 32 0 -%14 = OpConstant %13 0 -%15 = OpConstant %6 -1 -%16 = OpConstant %13 1 -%17 = OpConstant %6 0 -%18 = OpConstant %13 4294967295 -%19 = OpConstant %6 1 -%27 = OpConstantComposite %3 %11 %11 %11 %11 -%28 = OpConstantComposite %3 %10 %10 %10 %10 -%31 = OpConstantNull %6 -%44 = OpTypeVector %13 2 -%54 = OpConstant %13 32 -%64 = OpConstantComposite %44 %54 %54 -%76 = OpConstant %6 31 -%82 = OpConstantComposite %5 %76 %76 -%8 = OpFunction %2 None %9 -%7 = OpLabel -OpBranch %20 -%20 = OpLabel -%21 = OpCompositeConstruct %3 %11 %11 %11 %11 -%22 = OpExtInst %4 %1 Degrees %10 -%23 = OpExtInst %4 %1 Radians %10 -%24 = OpExtInst %3 %1 Degrees %21 -%25 = OpExtInst %3 %1 Radians %21 -%26 = OpExtInst %3 %1 FClamp %21 %27 %28 -%29 = OpExtInst %3 %1 Refract %21 %21 %10 -%32 = OpCompositeExtract %6 %12 0 -%33 = OpCompositeExtract %6 %12 0 -%34 = OpIMul %6 %32 %33 -%35 = OpIAdd %6 %31 %34 -%36 = OpCompositeExtract %6 %12 1 -%37 = OpCompositeExtract %6 %12 1 -%38 = OpIMul %6 %36 %37 -%30 = OpIAdd %6 %35 %38 -%39 = OpCopyObject %13 %14 -%40 = OpExtInst %13 %1 FindUMsb %39 -%41 = OpExtInst %6 %1 FindSMsb %15 -%42 = OpCompositeConstruct %5 %15 %15 -%43 = OpExtInst %5 %1 FindSMsb %42 -%45 = OpCompositeConstruct %44 %16 %16 -%46 = OpExtInst %44 %1 FindUMsb %45 -%47 = OpExtInst %6 %1 FindILsb %15 -%48 = OpExtInst %13 %1 FindILsb %16 -%49 = OpCompositeConstruct %5 %15 %15 -%50 = OpExtInst %5 %1 FindILsb %49 -%51 = OpCompositeConstruct %44 %16 %16 -%52 = OpExtInst %44 %1 FindILsb %51 -%55 = OpExtInst %13 %1 FindILsb %14 -%53 = OpExtInst %13 %1 UMin %54 %55 -%57 = OpExtInst %6 %1 FindILsb %17 -%56 = OpExtInst %6 %1 UMin %54 %57 -%59 = OpExtInst %13 %1 FindILsb %18 -%58 = OpExtInst %13 %1 UMin %54 %59 -%61 = OpExtInst %6 %1 FindILsb %15 -%60 = OpExtInst %6 %1 UMin %54 %61 -%62 = OpCompositeConstruct %44 %14 %14 -%65 = OpExtInst %44 %1 FindILsb %62 -%63 = OpExtInst %44 %1 UMin %64 %65 -%66 = OpCompositeConstruct %5 %17 %17 -%68 = OpExtInst %5 %1 FindILsb %66 -%67 = OpExtInst %5 %1 UMin %64 %68 -%69 = OpCompositeConstruct %44 %16 %16 -%71 = OpExtInst %44 %1 FindILsb %69 -%70 = OpExtInst %44 %1 UMin %64 %71 -%72 = OpCompositeConstruct %5 %19 %19 -%74 = OpExtInst %5 %1 FindILsb %72 -%73 = OpExtInst %5 %1 UMin %64 %74 -%77 = OpExtInst %6 %1 FindUMsb %15 -%75 = OpISub %6 %76 %77 -%79 = OpExtInst %6 %1 FindUMsb %16 -%78 = OpISub %13 %76 %79 -%80 = OpCompositeConstruct %5 %15 %15 -%83 = OpExtInst %5 %1 FindUMsb %80 -%81 = OpISub %5 %82 %83 -%84 = OpCompositeConstruct %44 %16 %16 -%86 = OpExtInst %5 %1 FindUMsb %84 -%85 = OpISub %44 %82 %86 +%7 = OpTypeStruct %4 %4 +%8 = OpTypeVector %4 2 +%9 = OpTypeStruct %8 %8 +%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 %4 1.5 +%35 = OpConstantComposite %3 %18 %18 %18 %18 +%36 = OpConstantComposite %3 %17 %17 %17 %17 +%39 = OpConstantNull %6 +%52 = OpTypeVector %20 2 +%62 = OpConstant %20 32 +%72 = OpConstantComposite %52 %62 %62 +%84 = OpConstant %6 31 +%90 = OpConstantComposite %5 %84 %84 +%15 = OpFunction %2 None %16 +%14 = OpLabel +OpBranch %28 +%28 = OpLabel +%29 = OpCompositeConstruct %3 %18 %18 %18 %18 +%30 = OpExtInst %4 %1 Degrees %17 +%31 = OpExtInst %4 %1 Radians %17 +%32 = OpExtInst %3 %1 Degrees %29 +%33 = OpExtInst %3 %1 Radians %29 +%34 = OpExtInst %3 %1 FClamp %29 %35 %36 +%37 = OpExtInst %3 %1 Refract %29 %29 %17 +%40 = OpCompositeExtract %6 %19 0 +%41 = OpCompositeExtract %6 %19 0 +%42 = OpIMul %6 %40 %41 +%43 = OpIAdd %6 %39 %42 +%44 = OpCompositeExtract %6 %19 1 +%45 = OpCompositeExtract %6 %19 1 +%46 = OpIMul %6 %44 %45 +%38 = OpIAdd %6 %43 %46 +%47 = OpCopyObject %20 %21 +%48 = OpExtInst %20 %1 FindUMsb %47 +%49 = OpExtInst %6 %1 FindSMsb %22 +%50 = OpCompositeConstruct %5 %22 %22 +%51 = OpExtInst %5 %1 FindSMsb %50 +%53 = OpCompositeConstruct %52 %23 %23 +%54 = OpExtInst %52 %1 FindUMsb %53 +%55 = OpExtInst %6 %1 FindILsb %22 +%56 = OpExtInst %20 %1 FindILsb %23 +%57 = OpCompositeConstruct %5 %22 %22 +%58 = OpExtInst %5 %1 FindILsb %57 +%59 = OpCompositeConstruct %52 %23 %23 +%60 = OpExtInst %52 %1 FindILsb %59 +%63 = OpExtInst %20 %1 FindILsb %21 +%61 = OpExtInst %20 %1 UMin %62 %63 +%65 = OpExtInst %6 %1 FindILsb %24 +%64 = OpExtInst %6 %1 UMin %62 %65 +%67 = OpExtInst %20 %1 FindILsb %25 +%66 = OpExtInst %20 %1 UMin %62 %67 +%69 = OpExtInst %6 %1 FindILsb %22 +%68 = OpExtInst %6 %1 UMin %62 %69 +%70 = OpCompositeConstruct %52 %21 %21 +%73 = OpExtInst %52 %1 FindILsb %70 +%71 = OpExtInst %52 %1 UMin %72 %73 +%74 = OpCompositeConstruct %5 %24 %24 +%76 = OpExtInst %5 %1 FindILsb %74 +%75 = OpExtInst %5 %1 UMin %72 %76 +%77 = OpCompositeConstruct %52 %23 %23 +%79 = OpExtInst %52 %1 FindILsb %77 +%78 = OpExtInst %52 %1 UMin %72 %79 +%80 = OpCompositeConstruct %5 %26 %26 +%82 = OpExtInst %5 %1 FindILsb %80 +%81 = OpExtInst %5 %1 UMin %72 %82 +%85 = OpExtInst %6 %1 FindUMsb %22 +%83 = OpISub %6 %84 %85 +%87 = OpExtInst %6 %1 FindUMsb %23 +%86 = OpISub %20 %84 %87 +%88 = OpCompositeConstruct %5 %22 %22 +%91 = OpExtInst %5 %1 FindUMsb %88 +%89 = OpISub %5 %90 %91 +%92 = OpCompositeConstruct %52 %23 %23 +%94 = OpExtInst %5 %1 FindUMsb %92 +%93 = OpISub %52 %90 %94 +%95 = OpExtInst %7 %1 ModfStruct %27 +%96 = OpExtInst %7 %1 ModfStruct %27 +%97 = OpCompositeExtract %4 %96 0 +%98 = OpExtInst %7 %1 ModfStruct %27 +%99 = OpCompositeExtract %4 %98 1 +%100 = OpCompositeConstruct %8 %27 %27 +%101 = OpExtInst %9 %1 ModfStruct %100 +%102 = OpCompositeConstruct %3 %27 %27 %27 %27 +%103 = OpExtInst %10 %1 ModfStruct %102 +%104 = OpCompositeExtract %3 %103 1 +%105 = OpCompositeExtract %4 %104 0 +%106 = OpCompositeConstruct %8 %27 %27 +%107 = OpExtInst %9 %1 ModfStruct %106 +%108 = OpCompositeExtract %8 %107 0 +%109 = OpCompositeExtract %4 %108 1 +%110 = OpExtInst %11 %1 FrexpStruct %27 +%111 = OpExtInst %11 %1 FrexpStruct %27 +%112 = OpCompositeExtract %4 %111 0 +%113 = OpExtInst %11 %1 FrexpStruct %27 +%114 = OpCompositeExtract %6 %113 1 +%115 = OpCompositeConstruct %3 %27 %27 %27 %27 +%116 = OpExtInst %13 %1 FrexpStruct %115 +%117 = OpCompositeExtract %12 %116 1 +%118 = OpCompositeExtract %6 %117 0 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/atomicCompareExchange.wgsl b/tests/out/wgsl/atomicCompareExchange.wgsl index 2ff428bd71..b659bed2b5 100644 --- a/tests/out/wgsl/atomicCompareExchange.wgsl +++ b/tests/out/wgsl/atomicCompareExchange.wgsl @@ -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) diff --git a/tests/out/wgsl/math-functions.wgsl b/tests/out/wgsl/math-functions.wgsl index 5faccba7b8..1eded4c04d 100644 --- a/tests/out/wgsl/math-functions.wgsl +++ b/tests/out/wgsl/math-functions.wgsl @@ -28,4 +28,14 @@ fn main() { let clz_b = countLeadingZeros(1u); let clz_c = countLeadingZeros(vec2(-1)); let clz_d = countLeadingZeros(vec2(1u)); + 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 = modf(vec2(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(1.5, 1.5, 1.5, 1.5)).exp.x; }