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 4b732a6038..54d5f341a2 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; @@ -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", 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 c37334b60d..6da98502f1 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 @@ -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"), 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 43d094ec2e..da39cd94f2 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -97,6 +97,14 @@ impl Writer { self.ep_results.clear(); } + fn is_builtin_wgsl_struct(&self, module: &Module, handle: Handle) -> 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 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(module, handle) { + 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..0c608504c9 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, 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 + } } diff --git a/src/front/wgsl/lower/mod.rs b/src/front/wgsl/lower/mod.rs index 23d3282364..64ef43cb93 100644 --- a/src/front/wgsl/lower/mod.rs +++ b/src/front/wgsl/lower/mod.rs @@ -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, } } 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 fdea215e33..49fd9d4e5a 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -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, + 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))] @@ -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>, + + /// 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 f0131c1df0..548b6c1451 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -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 { - 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::Ldexp => { diff --git a/tests/in/math-functions.wgsl b/tests/in/math-functions.wgsl index afc90811ca..408f8a74f8 100644 --- a/tests/in/math-functions.wgsl +++ b/tests/in/math-functions.wgsl @@ -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; } diff --git a/tests/out/glsl/math-functions.main.Fragment.glsl b/tests/out/glsl/math-functions.main.Fragment.glsl index 02cafcea02..be81715ce1 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); @@ -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; } diff --git a/tests/out/hlsl/math-functions.hlsl b/tests/out/hlsl/math-functions.hlsl index 3061a3ac1b..afdc6f4671 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; @@ -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; } diff --git a/tests/out/ir/access.ron b/tests/out/ir/access.ron index 5cda4104cd..a200e75a91 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, + predeclared_types: {}, ), constants: [], global_variables: [ diff --git a/tests/out/ir/collatz.ron b/tests/out/ir/collatz.ron index 6451fc715f..fb8ff825e1 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, + predeclared_types: {}, ), constants: [], global_variables: [ diff --git a/tests/out/ir/shadow.ron b/tests/out/ir/shadow.ron index 9ae5af0003..bfde76d6ae 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, + predeclared_types: {}, ), constants: [ ( diff --git a/tests/out/msl/math-functions.msl b/tests/out/msl/math-functions.msl index 1f22ae81dc..14824ec30f 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_( ) { @@ -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; } diff --git a/tests/out/spv/math-functions.spvasm b/tests/out/spv/math-functions.spvasm index b3a32bce33..260c3b4bd4 100644 --- a/tests/out/spv/math-functions.spvasm +++ b/tests/out/spv/math-functions.spvasm @@ -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 \ 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 5cdab8cdb7..149ebff8e0 100644 --- a/tests/out/wgsl/math-functions.wgsl +++ b/tests/out/wgsl/math-functions.wgsl @@ -30,4 +30,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 = 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; }