From 0a7b8ee15777932ee28aa26d7abeeec353625469 Mon Sep 17 00:00:00 2001 From: Fredrik Fornwall Date: Mon, 21 Aug 2023 10:19:57 +0200 Subject: [PATCH] [wgsl-in] Handle modf and frexp --- src/back/glsl/keywords.rs | 5 + src/back/glsl/mod.rs | 57 +++++- src/back/hlsl/help.rs | 48 +++++ src/back/hlsl/keywords.rs | 5 + src/back/hlsl/writer.rs | 25 ++- src/back/msl/keywords.rs | 4 + src/back/msl/writer.rs | 90 ++++++++- src/back/spv/block.rs | 15 +- src/back/spv/mod.rs | 2 + src/back/spv/writer.rs | 65 +++++- src/back/wgsl/writer.rs | 6 + src/front/type_gen.rs | 32 +++ src/front/wgsl/lower/mod.rs | 98 +++++++--- src/front/wgsl/mod.rs | 2 + src/front/wgsl/tests.rs | 5 +- src/lib.rs | 33 +++- src/proc/layouter.rs | 4 + src/proc/mod.rs | 5 +- src/proc/typifier.rs | 48 ++++- src/valid/expression.rs | 34 +--- src/valid/handles.rs | 2 + src/valid/mod.rs | 2 + src/valid/type.rs | 3 + tests/in/math-functions.wgsl | 6 + .../glsl/math-functions.main.Fragment.glsl | 26 +++ tests/out/hlsl/math-functions.hlsl | 34 ++++ tests/out/ir/access.ron | 2 + tests/out/ir/collatz.ron | 2 + tests/out/ir/shadow.ron | 2 + tests/out/msl/math-functions.msl | 28 +++ tests/out/spv/math-functions.spvasm | 185 ++++++++++-------- tests/out/wgsl/math-functions.wgsl | 6 + 32 files changed, 727 insertions(+), 154 deletions(-) diff --git a/src/back/glsl/keywords.rs b/src/back/glsl/keywords.rs index 9679020a0a..e6e78b643a 100644 --- a/src/back/glsl/keywords.rs +++ b/src/back/glsl/keywords.rs @@ -477,4 +477,9 @@ pub const RESERVED_KEYWORDS: &[&str] = &[ // entry point name (should not be shadowed) // "main", + // Naga utilities: + super::MODF_FUNCTION, + super::MODF_STRUCT, + super::FREXP_FUNCTION, + super::FREXP_STRUCT, ]; diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index ef5dd0143f..6a934c0a57 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -72,6 +72,11 @@ 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 FREXP_FUNCTION: &str = "naga_frexp"; +pub(crate) const FREXP_STRUCT: &str = "naga_frexp_struct"; +pub(crate) const MODF_FUNCTION: &str = "naga_modf"; +pub(crate) const MODF_STRUCT: &str = "naga_modf_struct"; + /// Mapping between resources and bindings. pub type BindingMap = std::collections::BTreeMap; @@ -604,6 +609,38 @@ impl<'a, W: Write> Writer<'a, W> { } } + if self.module.special_types.modf_result.is_some() { + writeln!( + self.out, + "struct {MODF_STRUCT} {{ + float fract; + float whole; +}}; + +{MODF_STRUCT} {MODF_FUNCTION}(float arg) {{ + float whole; + float fract = modf(arg, whole); + return {MODF_STRUCT}(fract, whole); +}}" + )?; + } + + if self.module.special_types.frexp_result.is_some() { + writeln!( + self.out, + "struct {FREXP_STRUCT} {{ + float fract; + int exp; +}}; + +{FREXP_STRUCT} {FREXP_FUNCTION}(float arg) {{ + int exp; + float fract = frexp(arg, exp); + return {FREXP_STRUCT}(fract, exp); +}}" + )?; + } + // Write struct types. // // This are always ordered because the IR is structured in a way that @@ -860,6 +897,8 @@ impl<'a, W: Write> Writer<'a, W> { | TypeInner::Sampler { .. } | TypeInner::AccelerationStructure | TypeInner::RayQuery + | TypeInner::ModfResult + | TypeInner::FrexpResult | TypeInner::BindingArray { .. } => { return Err(Error::Custom(format!("Unable to write type {inner:?}"))) } @@ -885,6 +924,14 @@ impl<'a, W: Write> Writer<'a, W> { } // glsl array has the size separated from the base type TypeInner::Array { base, .. } => self.write_type(base), + TypeInner::FrexpResult => { + write!(self.out, "{FREXP_STRUCT}")?; + Ok(()) + } + TypeInner::ModfResult => { + write!(self.out, "{MODF_STRUCT}")?; + Ok(()) + } ref other => self.write_value_type(other), } } @@ -2325,6 +2372,12 @@ impl<'a, W: Write> Writer<'a, W> { &self.names[&NameKey::StructMember(ty, index)] )? } + TypeInner::FrexpResult => { + write!(self.out, ".{}", if index == 0 { "fract" } else { "exp " })? + } + TypeInner::ModfResult => { + write!(self.out, ".{}", if index == 0 { "fract" } else { "whole " })? + } ref other => return Err(Error::Custom(format!("Cannot index {other:?}"))), } } @@ -2985,8 +3038,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..5129c4aae4 100644 --- a/src/back/hlsl/help.rs +++ b/src/back/hlsl/help.rs @@ -781,6 +781,54 @@ impl<'a, W: Write> super::Writer<'a, W> { Ok(()) } + pub(super) fn write_special_functions(&mut self, module: &crate::Module) -> BackendResult { + if module.special_types.frexp_result.is_some() { + let function_name = super::writer::FREXP_FUNCTION; + let struct_name = super::writer::FREXP_STRUCT; + writeln!( + self.out, + "struct {struct_name} {{ + float fract; + int exp; +}}; + +{struct_name} {function_name}(in float arg) {{ + float exp; + float fract = frexp(arg, exp); + {struct_name} result; + result.exp = exp; + result.fract = fract; + return result; +}}" + )?; + // Write extra new line + writeln!(self.out)?; + } + if module.special_types.modf_result.is_some() { + let function_name = super::writer::MODF_FUNCTION; + let struct_name = super::writer::MODF_STRUCT; + writeln!( + self.out, + "struct {struct_name} {{ + float fract; + float whole; +}}; + +{struct_name} {function_name}(in float arg) {{ + float whole; + float fract = modf(arg, whole); + {struct_name} result; + result.whole = whole; + result.fract = fract; + return result; +}}" + )?; + // Write extra new line + writeln!(self.out)?; + } + 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..b900c1446e 100644 --- a/src/back/hlsl/keywords.rs +++ b/src/back/hlsl/keywords.rs @@ -814,6 +814,11 @@ pub const RESERVED: &[&str] = &[ "TextureBuffer", "ConstantBuffer", "RayQuery", + // Naga utilities + super::writer::FREXP_FUNCTION, + super::writer::FREXP_STRUCT, + super::writer::MODF_FUNCTION, + super::writer::MODF_STRUCT, ]; // 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..62b044c6b8 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -17,6 +17,11 @@ const SPECIAL_BASE_VERTEX: &str = "base_vertex"; const SPECIAL_BASE_INSTANCE: &str = "base_instance"; const SPECIAL_OTHER: &str = "other"; +pub(crate) const FREXP_FUNCTION: &str = "frexp_modf"; +pub(crate) const FREXP_STRUCT: &str = "frexp_modf_result"; +pub(crate) const MODF_FUNCTION: &str = "naga_modf"; +pub(crate) const MODF_STRUCT: &str = "naga_modf_result"; + struct EpStructMember { name: String, ty: Handle, @@ -244,6 +249,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 @@ -1058,6 +1065,12 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { TypeInner::Array { base, size, .. } | TypeInner::BindingArray { base, size } => { self.write_array_size(module, base, size)?; } + TypeInner::FrexpResult => { + write!(self.out, "struct {FREXP_STRUCT}")?; + } + TypeInner::ModfResult => { + write!(self.out, "struct {MODF_STRUCT}")?; + } _ => return Err(Error::Unimplemented(format!("write_value_type {inner:?}"))), } @@ -2276,6 +2289,14 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { &writer.names[&NameKey::StructMember(ty, index)] )? } + TypeInner::FrexpResult => { + write!(writer.out, ".{}", if index == 0 { "fract" } else { "exp" })? + } + TypeInner::ModfResult => write!( + writer.out, + ".{}", + if index == 0 { "fract" } else { "whole" } + )?, ref other => { return Err(Error::Custom(format!("Cannot index {other:?}"))) } @@ -2665,8 +2686,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..27f86894af 100644 --- a/src/back/msl/keywords.rs +++ b/src/back/msl/keywords.rs @@ -214,4 +214,8 @@ pub const RESERVED: &[&str] = &[ // Naga utilities "DefaultConstructible", "clamped_lod_e", + super::writer::FREXP_FUNCTION, + super::writer::FREXP_STRUCT, + super::writer::MODF_FUNCTION, + super::writer::MODF_STRUCT, ]; diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index a8a103e6d0..bb0e3fb2cb 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -32,6 +32,11 @@ 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 FREXP_FUNCTION: &str = "naga_frexp"; +pub(crate) const FREXP_STRUCT: &str = "naga_frexp_result"; +pub(crate) const MODF_FUNCTION: &str = "naga_modf"; +pub(crate) const MODF_STRUCT: &str = "naga_modf_result"; + /// Write the Metal name for a Naga numeric type: scalar, vector, or matrix. /// /// The `sizes` slice determines whether this function writes a @@ -140,7 +145,15 @@ impl<'a> Display for TypeContext<'a> { // so just print the element type here. write!(out, "{sub}") } - crate::TypeInner::Struct { .. } => unreachable!(), + crate::TypeInner::Struct { .. } => { + unreachable!() + } + crate::TypeInner::FrexpResult { .. } => { + write!(out, "{}", FREXP_STRUCT) + } + crate::TypeInner::ModfResult { .. } => { + write!(out, "{}", MODF_STRUCT) + } crate::TypeInner::Image { dim, arrayed, @@ -452,6 +465,8 @@ impl crate::Type { | Ti::Sampler { .. } | Ti::AccelerationStructure | Ti::RayQuery + | Ti::FrexpResult + | Ti::ModfResult | Ti::BindingArray { .. } => false, } } @@ -1635,6 +1650,24 @@ impl Writer { write!(self.out, "{NAMESPACE}::{op}")?; self.put_call_parameters(iter::once(argument), context)?; } + crate::Expression::Math { + fun: crate::MathFunction::Frexp, + arg, + .. + } => { + write!(self.out, "{FREXP_FUNCTION}(")?; + self.put_expression(arg, context, false)?; + write!(self.out, ")")?; + } + crate::Expression::Math { + fun: crate::MathFunction::Modf, + arg, + .. + } => { + write!(self.out, "{MODF_FUNCTION}(")?; + self.put_expression(arg, context, false)?; + write!(self.out, ")")?; + } crate::Expression::Math { fun, arg, @@ -1644,7 +1677,7 @@ impl Writer { } => { use crate::MathFunction as Mf; - let scalar_argument = match *context.resolve_type(arg) { + let scalar_argument: bool = match *context.resolve_type(arg) { crate::TypeInner::Scalar { .. } => true, _ => false, }; @@ -2018,7 +2051,9 @@ impl Writer { base_inner = &context.module.types[base].inner; } match *base_inner { - crate::TypeInner::Struct { .. } => (base, None), + crate::TypeInner::Struct { .. } + | crate::TypeInner::FrexpResult + | crate::TypeInner::ModfResult { .. } => (base, None), _ => (base, Some(index::GuardedIndex::Known(index))), } } @@ -2133,6 +2168,20 @@ impl Writer { write!(self.out, ".{}", back::COMPONENTS[index as usize])?; } } + crate::TypeInner::FrexpResult | crate::TypeInner::ModfResult => { + self.put_access_chain(base, policy, context)?; + write!( + self.out, + ".{}", + if index == 0 { + "fract" + } else if *base_ty == crate::TypeInner::FrexpResult { + "exp" + } else { + "whole" + } + )?; + } _ => { self.put_subscripted_access_chain( base, @@ -3236,6 +3285,41 @@ impl Writer { } } } + + if module.special_types.frexp_result.is_some() { + writeln!( + self.out, + " +struct {FREXP_STRUCT} {{ + float fract; + int exp; +}}; + +struct {FREXP_STRUCT} {FREXP_FUNCTION}(float arg) {{ + int exp; + float fract = {NAMESPACE}::frexp(arg, exp); + return {FREXP_STRUCT}{{ fract, exp }}; +}};" + )?; + } + + if module.special_types.modf_result.is_some() { + writeln!( + self.out, + " +struct {MODF_STRUCT} {{ + float fract; + float whole; +}}; + +struct {MODF_STRUCT} {MODF_FUNCTION}(float arg) {{ + float whole; + float fract = {NAMESPACE}::modf(arg, whole); + return {MODF_STRUCT}{{ fract, whole }}; +}};" + )?; + } + Ok(()) } diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index 11841bc522..2b8295498f 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -396,6 +396,17 @@ impl<'w> BlockContext<'w> { load_id } + crate::TypeInner::FrexpResult | crate::TypeInner::ModfResult => { + let id = self.gen_id(); + let base_id = self.cached[base]; + block.body.push(Instruction::composite_extract( + result_type_id, + id, + base_id, + &[index], + )); + id + } ref other => { log::error!("Unable to access index of {:?}", other); return Err(Error::FeatureNotImplemented("access index for type")); @@ -787,8 +798,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/spv/mod.rs b/src/back/spv/mod.rs index 52e304a237..ceb169aae9 100644 --- a/src/back/spv/mod.rs +++ b/src/back/spv/mod.rs @@ -403,6 +403,8 @@ fn make_local(inner: &crate::TypeInner) -> Option { crate::TypeInner::RayQuery => LocalType::RayQuery, crate::TypeInner::Array { .. } | crate::TypeInner::Struct { .. } + | crate::TypeInner::FrexpResult + | crate::TypeInner::ModfResult | crate::TypeInner::BindingArray { .. } => return None, }) } diff --git a/src/back/spv/writer.rs b/src/back/spv/writer.rs index e6db93602a..2733cc9ab8 100644 --- a/src/back/spv/writer.rs +++ b/src/back/spv/writer.rs @@ -241,6 +241,16 @@ impl Writer { self.get_type_id(local_type.into()) } + pub(super) fn get_int_type_id(&mut self) -> Word { + let local_type = LocalType::Value { + vector_size: None, + kind: crate::ScalarKind::Sint, + width: 4, + pointer_space: None, + }; + self.get_type_id(local_type.into()) + } + pub(super) fn get_float_type_id(&mut self) -> Word { let local_type = LocalType::Value { vector_size: None, @@ -1024,6 +1034,36 @@ impl Writer { } Instruction::type_struct(id, member_ids.as_slice()) } + crate::TypeInner::FrexpResult | crate::TypeInner::ModfResult => { + let ty_inner = crate::TypeInner::Scalar { + width: 4, + kind: crate::ScalarKind::Float, + }; + self.decorate_struct_member_raw(id, 0, 0, Some("fract"), &ty_inner, arena)?; + self.decorate_struct_member_raw( + id, + 1, + 4, + Some(if ty.inner == crate::TypeInner::FrexpResult { + "exp" + } else { + "whole" + }), + &ty_inner, + arena, + )?; + + let float_id = self.get_float_type_id(); + let member_ids = [ + float_id, + if ty.inner == crate::TypeInner::FrexpResult { + self.get_int_type_id() + } else { + float_id + }, + ]; + Instruction::type_struct(id, &member_ids) + } // These all have TypeLocal representations, so they should have been // handled by `write_type_declaration_local` above. @@ -1693,6 +1733,25 @@ impl Writer { index: usize, member: &crate::StructMember, arena: &UniqueArena, + ) -> Result<(), Error> { + self.decorate_struct_member_raw( + struct_id, + index, + member.offset, + member.name.as_deref(), + &arena[member.ty].inner, + arena, + ) + } + + fn decorate_struct_member_raw( + &mut self, + struct_id: Word, + index: usize, + offset: Word, + member_name: Option<&str>, + member_ty_inner: &crate::TypeInner, + arena: &UniqueArena, ) -> Result<(), Error> { use spirv::Decoration; @@ -1700,11 +1759,11 @@ impl Writer { struct_id, index as u32, Decoration::Offset, - &[member.offset], + &[offset], )); if self.flags.contains(WriterFlags::DEBUG) { - if let Some(ref name) = member.name { + if let Some(name) = member_name { self.debugs .push(Instruction::member_name(struct_id, index as u32, name)); } @@ -1712,7 +1771,7 @@ impl Writer { // Matrices and arrays of matrices both require decorations, // so "see through" an array to determine if they're needed. - let member_array_subty_inner = match arena[member.ty].inner { + let member_array_subty_inner = match *member_ty_inner { crate::TypeInner::Array { base, .. } => &arena[base].inner, ref other => other, }; diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index 44dd2a97c0..eb54acd3f1 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -1223,6 +1223,12 @@ impl Writer { &self.names[&NameKey::StructMember(ty, index)] )? } + TypeInner::FrexpResult => { + write!(self.out, ".{}", if index == 0 { "fract" } else { "exp" })? + } + TypeInner::ModfResult => { + write!(self.out, ".{}", if index == 0 { "fract" } else { "whole" })? + } ref other => return Err(Error::Custom(format!("Cannot index {other:?}"))), } } diff --git a/src/front/type_gen.rs b/src/front/type_gen.rs index 1ee454c448..7b3033aa4d 100644 --- a/src/front/type_gen.rs +++ b/src/front/type_gen.rs @@ -311,4 +311,36 @@ impl crate::Module { self.special_types.ray_intersection = Some(handle); handle } + + /// Populate this module's [`crate::SpecialTypes::modf_result`] type. + pub fn generate_modf_result(&mut self) -> Handle { + if let Some(handle) = self.special_types.modf_result { + return handle; + } + let handle = self.types.insert( + crate::Type { + name: Some("__naga_modf_result".to_string()), + inner: crate::TypeInner::ModfResult, + }, + Span::UNDEFINED, + ); + self.special_types.modf_result = Some(handle); + handle + } + + /// Populate this module's [`crate::SpecialTypes::frexp_result`] type. + pub fn generate_frexp_result(&mut self) -> Handle { + if let Some(handle) = self.special_types.frexp_result { + return handle; + } + let handle = self.types.insert( + crate::Type { + name: Some("__naga_frexp_result".to_string()), + inner: crate::TypeInner::FrexpResult, + }, + Span::UNDEFINED, + ); + self.special_types.frexp_result = Some(handle); + handle + } } diff --git a/src/front/wgsl/lower/mod.rs b/src/front/wgsl/lower/mod.rs index f3a7b22dd3..80791f1f28 100644 --- a/src/front/wgsl/lower/mod.rs +++ b/src/front/wgsl/lower/mod.rs @@ -1587,6 +1587,38 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { ), } } + crate::TypeInner::FrexpResult => { + let index = if field.name == "fract" { + 0 + } else if field.name == "exp" { + 1 + } else { + return Err(Error::BadAccessor(field.span)); + }; + ( + crate::Expression::AccessIndex { + base: handle, + index, + }, + is_reference, + ) + } + crate::TypeInner::ModfResult => { + let index = if field.name == "fract" { + 0 + } else if field.name == "whole" { + 1 + } else { + return Err(Error::BadAccessor(field.span)); + }; + ( + crate::Expression::AccessIndex { + base: handle, + index, + }, + is_reference, + ) + } _ => return Err(Error::BadAccessor(field.span)), }; @@ -1711,32 +1743,52 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { } else if let Some(fun) = conv::map_standard_fun(function.name) { let expected = fun.argument_count() as _; let mut args = ctx.prepare_args(arguments, expected, span); - let arg = self.expression(args.next()?, ctx.reborrow())?; - let arg1 = args - .next() - .map(|x| self.expression(x, ctx.reborrow())) - .ok() - .transpose()?; - let arg2 = args - .next() - .map(|x| self.expression(x, ctx.reborrow())) - .ok() - .transpose()?; - let arg3 = args - .next() - .map(|x| self.expression(x, ctx.reborrow())) - .ok() - .transpose()?; - args.finish()?; + // Handle special functions return structures + match fun { + crate::MathFunction::Frexp | crate::MathFunction::Modf => { + args.finish()?; + let _ = if fun == crate::MathFunction::Frexp { + ctx.module.generate_frexp_result() + } else { + ctx.module.generate_modf_result() + }; + crate::Expression::Math { + fun, + arg, + arg1: None, + arg2: None, + arg3: None, + } + } + _ => { + let arg1 = args + .next() + .map(|x| self.expression(x, ctx.reborrow())) + .ok() + .transpose()?; + let arg2 = args + .next() + .map(|x| self.expression(x, ctx.reborrow())) + .ok() + .transpose()?; + let arg3 = args + .next() + .map(|x| self.expression(x, ctx.reborrow())) + .ok() + .transpose()?; - crate::Expression::Math { - fun, - arg, - arg1, - arg2, - arg3, + args.finish()?; + + crate::Expression::Math { + fun, + arg, + arg1, + arg2, + arg3, + } + } } } else if let Some(fun) = Texture::map(function.name) { self.texture_sample_helper(fun, arguments, span, ctx.reborrow())? diff --git a/src/front/wgsl/mod.rs b/src/front/wgsl/mod.rs index 2b471a8a93..c7ff07559d 100644 --- a/src/front/wgsl/mod.rs +++ b/src/front/wgsl/mod.rs @@ -185,6 +185,8 @@ impl crate::TypeInner { Ti::Sampler { .. } => "sampler".to_string(), Ti::AccelerationStructure => "acceleration_structure".to_string(), Ti::RayQuery => "ray_query".to_string(), + Ti::FrexpResult => "__frexp_result_f32".to_string(), + Ti::ModfResult => "__modf_result_f32".to_string(), Ti::BindingArray { base, size, .. } => { let member_type = &gctx.types[base]; let base = member_type.name.as_deref().unwrap_or("unknown"); diff --git a/src/front/wgsl/tests.rs b/src/front/wgsl/tests.rs index 02fc110cae..3f39c7cb44 100644 --- a/src/front/wgsl/tests.rs +++ b/src/front/wgsl/tests.rs @@ -438,10 +438,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..242bc0f91f 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -672,7 +672,10 @@ pub struct Type { #[cfg_attr(feature = "arbitrary", derive(Arbitrary))] pub enum TypeInner { /// Number of integral or floating-point kind. - Scalar { kind: ScalarKind, width: Bytes }, + Scalar { + kind: ScalarKind, + width: Bytes, + }, /// Vector of numbers. Vector { size: VectorSize, @@ -686,7 +689,10 @@ pub enum TypeInner { width: Bytes, }, /// Atomic scalar. - Atomic { kind: ScalarKind, width: Bytes }, + Atomic { + kind: ScalarKind, + width: Bytes, + }, /// Pointer to another type. /// /// Pointers to scalars and vectors should be treated as equivalent to @@ -795,7 +801,9 @@ pub enum TypeInner { class: ImageClass, }, /// Can be used to sample values from images. - Sampler { comparison: bool }, + Sampler { + comparison: bool, + }, /// Opaque object representing an acceleration structure of geometry. AccelerationStructure, @@ -803,6 +811,12 @@ pub enum TypeInner { /// Locally used handle for ray queries. RayQuery, + // TODO: document + FrexpResult, + + // TODO: document + ModfResult, + /// Array of bindings. /// /// A `BindingArray` represents an array where each element draws its value @@ -842,7 +856,10 @@ pub enum TypeInner { /// [`DATA`]: crate::valid::TypeFlags::DATA /// [`ARGUMENT`]: crate::valid::TypeFlags::ARGUMENT /// [naga#1864]: https://github.com/gfx-rs/naga/issues/1864 - BindingArray { base: Handle, size: ArraySize }, + BindingArray { + base: Handle, + size: ArraySize, + }, } #[derive(Debug, Clone, Copy, PartialOrd)] @@ -1961,6 +1978,14 @@ pub struct SpecialTypes { /// Call [`Module::generate_ray_intersection_type`] to populate /// this if needed and return the handle. pub ray_intersection: Option>, + + // Type for `__modf_result_f32`. + // + pub modf_result: Option>, + + // Type for `__frexp_result_f32 `. + // + pub frexp_result: Option>, } /// Shader module. diff --git a/src/proc/layouter.rs b/src/proc/layouter.rs index 11b2250e93..7ee902aa9c 100644 --- a/src/proc/layouter.rs +++ b/src/proc/layouter.rs @@ -234,6 +234,10 @@ impl Layouter { alignment, } } + Ti::FrexpResult { .. } | Ti::ModfResult { .. } => TypeLayout { + size: 8, + alignment: Alignment::EIGHT, + }, Ti::Image { .. } | Ti::Sampler { .. } | Ti::AccelerationStructure diff --git a/src/proc/mod.rs b/src/proc/mod.rs index 582c887034..60dd7f06ad 100644 --- a/src/proc/mod.rs +++ b/src/proc/mod.rs @@ -229,6 +229,7 @@ impl super::TypeInner { | Self::AccelerationStructure | Self::RayQuery | Self::BindingArray { .. } => 0, + Self::FrexpResult { .. } | Self::ModfResult { .. } => 8, } } @@ -375,8 +376,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..4b3e4cd489 100644 --- a/src/proc/typifier.rs +++ b/src/proc/typifier.rs @@ -391,6 +391,22 @@ impl<'a> ResolveContext<'a> { } } Ti::BindingArray { base, .. } => Ti::Pointer { base, space }, + Ti::FrexpResult if (0..=1).contains(&index) => Ti::ValuePointer { + size: None, + kind: if index == 0 { + crate::ScalarKind::Float + } else { + crate::ScalarKind::Sint + }, + width: 4, + space, + }, + Ti::ModfResult if (0..=1).contains(&index) => Ti::ValuePointer { + size: None, + kind: crate::ScalarKind::Float, + width: 4, + space, + }, ref other => { log::error!("Access index sub-type {:?}", other); return Err(ResolveError::InvalidSubAccess { @@ -400,6 +416,22 @@ impl<'a> ResolveContext<'a> { } }), Ti::BindingArray { base, .. } => TypeResolution::Handle(base), + Ti::FrexpResult if (0..=1).contains(&index) => { + TypeResolution::Value(Ti::Scalar { + kind: if index == 0 { + crate::ScalarKind::Float + } else { + crate::ScalarKind::Sint + }, + width: 4, + }) + } + Ti::ModfResult if (0..=1).contains(&index) => { + TypeResolution::Value(Ti::Scalar { + kind: crate::ScalarKind::Float, + width: 4, + }) + } ref other => { log::error!("Access index type {:?}", other); return Err(ResolveError::InvalidAccess { @@ -706,8 +738,6 @@ impl<'a> ResolveContext<'a> { Mf::Round | Mf::Fract | Mf::Trunc | - Mf::Modf | - Mf::Frexp | Mf::Ldexp | // exponent Mf::Exp | @@ -715,6 +745,20 @@ impl<'a> ResolveContext<'a> { Mf::Log | Mf::Log2 | Mf::Pow => res_arg.clone(), + Mf::Frexp => { + let result = self + .special_types + .frexp_result + .ok_or(ResolveError::MissingSpecialType)?; + TypeResolution::Handle(result) + }, + Mf::Modf => { + let result = self + .special_types + .modf_result + .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..164b68389c 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -270,6 +270,7 @@ impl super::Validator { resolve_index_limit(module, top, &module.types[base].inner, false)? } Ti::Struct { ref members, .. } => members.len() as u32, + Ti::FrexpResult | Ti::ModfResult => 2, ref other => { log::error!("Indexing of {:?}", other); return Err(ExpressionError::InvalidBaseType(top)); @@ -1015,33 +1016,16 @@ 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, - kind: Sk::Float, - width, - space: _, - } => size == size0 && width == width0, - _ => false, - }; - if !good { + width: 4, + }, + ) { return Err(ExpressionError::InvalidArgumentType( fun, 1, diff --git a/src/valid/handles.rs b/src/valid/handles.rs index da95f60842..e0a8450931 100644 --- a/src/valid/handles.rs +++ b/src/valid/handles.rs @@ -54,6 +54,8 @@ impl super::Validator { | crate::TypeInner::Image { .. } | crate::TypeInner::Sampler { .. } | crate::TypeInner::AccelerationStructure + | crate::TypeInner::FrexpResult + | crate::TypeInner::ModfResult | crate::TypeInner::RayQuery => (), crate::TypeInner::Pointer { base, space: _ } => { this_handle.check_dep(base)?; diff --git a/src/valid/mod.rs b/src/valid/mod.rs index b04d69b5b4..be0fd6a31f 100644 --- a/src/valid/mod.rs +++ b/src/valid/mod.rs @@ -255,6 +255,8 @@ impl crate::TypeInner { | Self::Atomic { .. } | Self::Pointer { .. } | Self::ValuePointer { .. } + | Self::FrexpResult { .. } + | Self::ModfResult { .. } | Self::Struct { .. } => true, Self::Array { .. } | Self::Image { .. } diff --git a/src/valid/type.rs b/src/valid/type.rs index f8ceb463c6..0b423cff37 100644 --- a/src/valid/type.rs +++ b/src/valid/type.rs @@ -568,6 +568,9 @@ impl super::Validator { self.require_type_capability(Capabilities::RAY_QUERY)?; TypeInfo::new(TypeFlags::DATA | TypeFlags::SIZED, Alignment::ONE) } + Ti::FrexpResult | Ti::ModfResult => { + TypeInfo::new(TypeFlags::DATA | TypeFlags::SIZED, Alignment::EIGHT) + } Ti::BindingArray { base, size } => { if base >= handle { return Err(TypeError::InvalidArrayBaseType(base)); diff --git a/tests/in/math-functions.wgsl b/tests/in/math-functions.wgsl index efeb988f5a..f7a4541939 100644 --- a/tests/in/math-functions.wgsl +++ b/tests/in/math-functions.wgsl @@ -29,4 +29,10 @@ 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 frexp_a = frexp(1.5); + let frexp_b = frexp(1.5).fract; + let frexp_c = frexp(1.5).exp; } diff --git a/tests/out/glsl/math-functions.main.Fragment.glsl b/tests/out/glsl/math-functions.main.Fragment.glsl index 70ad0b199a..32446531c5 100644 --- a/tests/out/glsl/math-functions.main.Fragment.glsl +++ b/tests/out/glsl/math-functions.main.Fragment.glsl @@ -3,6 +3,26 @@ precision highp float; precision highp int; +struct naga_modf_struct { + float fract; + float whole; +}; + +naga_modf_struct naga_modf(float arg) { + float whole; + float fract = modf(arg, whole); + return naga_modf_struct(fract, whole); +} +struct naga_frexp_struct { + float fract; + int exp; +}; + +naga_frexp_struct naga_frexp(float arg) { + int exp; + float fract = frexp(arg, exp); + return naga_frexp_struct(fract, exp); +} void main() { vec4 v = vec4(0.0); @@ -34,5 +54,11 @@ 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))); + naga_modf_struct modf_a = naga_modf(1.5); + float modf_b = naga_modf(1.5).fract; + float modf_c = naga_modf(1.5).whole ; + naga_frexp_struct frexp_a = naga_frexp(1.5); + float frexp_b = naga_frexp(1.5).fract; + int frexp_c = naga_frexp(1.5).exp ; } diff --git a/tests/out/hlsl/math-functions.hlsl b/tests/out/hlsl/math-functions.hlsl index a1be810532..7a2e912782 100644 --- a/tests/out/hlsl/math-functions.hlsl +++ b/tests/out/hlsl/math-functions.hlsl @@ -1,3 +1,31 @@ +struct frexp_modf_result { + float fract; + int exp; +}; + +frexp_modf_result frexp_modf(in float arg) { + float exp; + float fract = frexp(arg, exp); + frexp_modf_result result; + result.exp = exp; + result.fract = fract; + return result; +} + +struct naga_modf_result { + float fract; + float whole; +}; + +naga_modf_result naga_modf(in float arg) { + float whole; + float fract = modf(arg, whole); + naga_modf_result result; + result.whole = whole; + result.fract = fract; + return result; +} + void main() { float4 v = (0.0).xxxx; @@ -29,4 +57,10 @@ 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)); + struct naga_modf_result modf_a = naga_modf(1.5); + float modf_b = naga_modf(1.5).fract; + float modf_c = naga_modf(1.5).whole; + struct frexp_modf_result frexp_a = frexp_modf(1.5); + float frexp_b = frexp_modf(1.5).fract; + int frexp_c = frexp_modf(1.5).exp; } diff --git a/tests/out/ir/access.ron b/tests/out/ir/access.ron index 7447249127..4f272d1902 100644 --- a/tests/out/ir/access.ron +++ b/tests/out/ir/access.ron @@ -334,6 +334,8 @@ special_types: ( ray_desc: None, ray_intersection: None, + modf_result: None, + frexp_result: None, ), constants: [], global_variables: [ diff --git a/tests/out/ir/collatz.ron b/tests/out/ir/collatz.ron index 6451fc715f..012f0a7c93 100644 --- a/tests/out/ir/collatz.ron +++ b/tests/out/ir/collatz.ron @@ -41,6 +41,8 @@ special_types: ( ray_desc: None, ray_intersection: None, + modf_result: None, + frexp_result: None, ), constants: [], global_variables: [ diff --git a/tests/out/ir/shadow.ron b/tests/out/ir/shadow.ron index 246657255b..86efcaeda8 100644 --- a/tests/out/ir/shadow.ron +++ b/tests/out/ir/shadow.ron @@ -266,6 +266,8 @@ special_types: ( ray_desc: None, ray_intersection: None, + modf_result: None, + frexp_result: None, ), constants: [ ( diff --git a/tests/out/msl/math-functions.msl b/tests/out/msl/math-functions.msl index 120d526d44..961c0ea588 100644 --- a/tests/out/msl/math-functions.msl +++ b/tests/out/msl/math-functions.msl @@ -5,6 +5,28 @@ using metal::uint; +struct naga_frexp_result { + float fract; + int exp; +}; + +struct naga_frexp_result naga_frexp(float arg) { + int exp; + float fract = metal::frexp(arg, exp); + return naga_frexp_result{ fract, exp }; +}; + +struct naga_modf_result { + float fract; + float whole; +}; + +struct naga_modf_result naga_modf(float arg) { + float whole; + float fract = metal::modf(arg, whole); + return naga_modf_result{ fract, whole }; +}; + fragment void main_( ) { metal::float4 v = metal::float4(0.0); @@ -38,4 +60,10 @@ 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)); + naga_modf_result modf_a = naga_modf(1.5); + float modf_b = naga_modf(1.5).fract; + float modf_c = naga_modf(1.5).whole; + naga_frexp_result frexp_a = naga_frexp(1.5); + float frexp_b = naga_frexp(1.5).fract; + int frexp_c = naga_frexp(1.5).exp; } diff --git a/tests/out/spv/math-functions.spvasm b/tests/out/spv/math-functions.spvasm index 60bdd7ae0d..0aa6b17af2 100644 --- a/tests/out/spv/math-functions.spvasm +++ b/tests/out/spv/math-functions.spvasm @@ -1,97 +1,114 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 87 +; Bound: 100 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Fragment %8 "main" -OpExecutionMode %8 OriginUpperLeft +OpEntryPoint Fragment %10 "main" +OpExecutionMode %10 OriginUpperLeft +OpMemberDecorate %7 0 Offset 0 +OpMemberDecorate %7 1 Offset 4 +OpMemberDecorate %8 0 Offset 0 +OpMemberDecorate %8 1 Offset 4 %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 = OpTypeStruct %4 %6 +%11 = OpTypeFunction %2 +%12 = OpConstant %4 1.0 +%13 = OpConstant %4 0.0 +%14 = OpConstantNull %5 +%15 = OpTypeInt 32 0 +%16 = OpConstant %15 0 +%17 = OpConstant %6 -1 +%18 = OpConstant %15 1 +%19 = OpConstant %6 0 +%20 = OpConstant %15 4294967295 +%21 = OpConstant %6 1 +%22 = OpConstant %4 1.5 +%30 = OpConstantComposite %3 %13 %13 %13 %13 +%31 = OpConstantComposite %3 %12 %12 %12 %12 +%34 = OpConstantNull %6 +%47 = OpTypeVector %15 2 +%57 = OpConstant %15 32 +%67 = OpConstantComposite %47 %57 %57 +%79 = OpConstant %6 31 +%85 = OpConstantComposite %5 %79 %79 +%10 = OpFunction %2 None %11 +%9 = OpLabel +OpBranch %23 +%23 = OpLabel +%24 = OpCompositeConstruct %3 %13 %13 %13 %13 +%25 = OpExtInst %4 %1 Degrees %12 +%26 = OpExtInst %4 %1 Radians %12 +%27 = OpExtInst %3 %1 Degrees %24 +%28 = OpExtInst %3 %1 Radians %24 +%29 = OpExtInst %3 %1 FClamp %24 %30 %31 +%32 = OpExtInst %3 %1 Refract %24 %24 %12 +%35 = OpCompositeExtract %6 %14 0 +%36 = OpCompositeExtract %6 %14 0 +%37 = OpIMul %6 %35 %36 +%38 = OpIAdd %6 %34 %37 +%39 = OpCompositeExtract %6 %14 1 +%40 = OpCompositeExtract %6 %14 1 +%41 = OpIMul %6 %39 %40 +%33 = OpIAdd %6 %38 %41 +%42 = OpCopyObject %15 %16 +%43 = OpExtInst %15 %1 FindUMsb %42 +%44 = OpExtInst %6 %1 FindSMsb %17 +%45 = OpCompositeConstruct %5 %17 %17 +%46 = OpExtInst %5 %1 FindSMsb %45 +%48 = OpCompositeConstruct %47 %18 %18 +%49 = OpExtInst %47 %1 FindUMsb %48 +%50 = OpExtInst %6 %1 FindILsb %17 +%51 = OpExtInst %15 %1 FindILsb %18 +%52 = OpCompositeConstruct %5 %17 %17 +%53 = OpExtInst %5 %1 FindILsb %52 +%54 = OpCompositeConstruct %47 %18 %18 +%55 = OpExtInst %47 %1 FindILsb %54 +%58 = OpExtInst %15 %1 FindILsb %16 +%56 = OpExtInst %15 %1 UMin %57 %58 +%60 = OpExtInst %6 %1 FindILsb %19 +%59 = OpExtInst %6 %1 UMin %57 %60 +%62 = OpExtInst %15 %1 FindILsb %20 +%61 = OpExtInst %15 %1 UMin %57 %62 +%64 = OpExtInst %6 %1 FindILsb %17 +%63 = OpExtInst %6 %1 UMin %57 %64 +%65 = OpCompositeConstruct %47 %16 %16 +%68 = OpExtInst %47 %1 FindILsb %65 +%66 = OpExtInst %47 %1 UMin %67 %68 +%69 = OpCompositeConstruct %5 %19 %19 +%71 = OpExtInst %5 %1 FindILsb %69 +%70 = OpExtInst %5 %1 UMin %67 %71 +%72 = OpCompositeConstruct %47 %18 %18 +%74 = OpExtInst %47 %1 FindILsb %72 +%73 = OpExtInst %47 %1 UMin %67 %74 +%75 = OpCompositeConstruct %5 %21 %21 +%77 = OpExtInst %5 %1 FindILsb %75 +%76 = OpExtInst %5 %1 UMin %67 %77 +%80 = OpExtInst %6 %1 FindUMsb %17 +%78 = OpISub %6 %79 %80 +%82 = OpExtInst %6 %1 FindUMsb %18 +%81 = OpISub %15 %79 %82 +%83 = OpCompositeConstruct %5 %17 %17 +%86 = OpExtInst %5 %1 FindUMsb %83 +%84 = OpISub %5 %85 %86 +%87 = OpCompositeConstruct %47 %18 %18 +%89 = OpExtInst %5 %1 FindUMsb %87 +%88 = OpISub %47 %85 %89 +%90 = OpExtInst %7 %1 ModfStruct %22 +%91 = OpExtInst %7 %1 ModfStruct %22 +%92 = OpCompositeExtract %4 %91 0 +%93 = OpExtInst %7 %1 ModfStruct %22 +%94 = OpCompositeExtract %4 %93 1 +%95 = OpExtInst %8 %1 FrexpStruct %22 +%96 = OpExtInst %8 %1 FrexpStruct %22 +%97 = OpCompositeExtract %4 %96 0 +%98 = OpExtInst %8 %1 FrexpStruct %22 +%99 = OpCompositeExtract %6 %98 1 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/math-functions.wgsl b/tests/out/wgsl/math-functions.wgsl index 5faccba7b8..99b71ad745 100644 --- a/tests/out/wgsl/math-functions.wgsl +++ b/tests/out/wgsl/math-functions.wgsl @@ -28,4 +28,10 @@ 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 frexp_a = frexp(1.5); + let frexp_b = frexp(1.5).fract; + let frexp_c = frexp(1.5).exp; }