From 16e99df2c9b07dd6a3f904c568d8d6db97f90691 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 | 35 +++- src/back/hlsl/help.rs | 36 ++++ src/back/hlsl/keywords.rs | 5 + src/back/hlsl/writer.rs | 14 +- src/back/msl/keywords.rs | 4 + src/back/msl/writer.rs | 38 +++- src/back/spv/block.rs | 4 +- src/back/spv/writer.rs | 35 +++- src/front/mod.rs | 2 +- src/front/type_gen.rs | 100 ++++++++++ src/front/wgsl/lower/mod.rs | 7 + src/front/wgsl/tests.rs | 5 +- src/lib.rs | 8 + src/proc/mod.rs | 4 +- src/proc/typifier.rs | 16 +- src/valid/expression.rs | 33 +--- 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 | 24 +++ tests/out/spv/math-functions.spvasm | 185 ++++++++++-------- tests/out/wgsl/math-functions.wgsl | 16 ++ 26 files changed, 520 insertions(+), 128 deletions(-) diff --git a/src/back/glsl/keywords.rs b/src/back/glsl/keywords.rs index 9679020a0a..ea6f171fc0 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::FREXP_FUNCTION, + crate::front::type_gen::FREXP_RESULT_F32_STRUCT_NAME, + super::MODF_FUNCTION, + crate::front::type_gen::MODF_RESULT_F32_STRUCT_NAME, ]; diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index ef5dd0143f..72110466fd 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 FREXP_FUNCTION: &str = "naga_frexp"; +pub(crate) const MODF_FUNCTION: &str = "naga_modf"; + /// Mapping between resources and bindings. pub type BindingMap = std::collections::BTreeMap; @@ -625,6 +628,34 @@ impl<'a, W: Write> Writer<'a, W> { } } + // Write special + if let Some(ty_struct) = self.module.special_types.frexp_result { + let struct_name = &self.names[&NameKey::Type(ty_struct)]; + writeln!(self.out)?; + writeln!( + self.out, + "{struct_name} {FREXP_FUNCTION}(float arg) {{ + int exp; + float fract = frexp(arg, exp); + return {struct_name}(fract, exp); +}}" + )?; + } + + if let Some(ty_struct) = self.module.special_types.modf_result { + let struct_name = &self.names[&NameKey::Type(ty_struct)]; + writeln!(self.out)?; + writeln!( + self.out, + "{} {MODF_FUNCTION}(float arg) {{ + float whole; + float fract = modf(arg, whole); + return {}(fract, whole); +}}", + struct_name, struct_name + )?; + } + // Write all named constants let mut constants = self .module @@ -2985,8 +3016,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..85f7ec85f7 100644 --- a/src/back/hlsl/help.rs +++ b/src/back/hlsl/help.rs @@ -781,6 +781,42 @@ impl<'a, W: Write> super::Writer<'a, W> { Ok(()) } + pub(super) fn write_special_functions(&mut self, module: &crate::Module) -> BackendResult { + if let Some(ty_struct) = module.special_types.frexp_result { + let struct_name = &self.names[&NameKey::Type(ty_struct)]; + writeln!( + self.out, + "{struct_name} {}(in float arg) {{ + float exp; + float fract = frexp(arg, exp); + {struct_name} result; + result.exp = exp; + result.fract = fract; + return result; +}}", + super::writer::FREXP_FUNCTION + )?; + writeln!(self.out)?; + } + if let Some(ty_struct) = module.special_types.modf_result { + let struct_name = &self.names[&NameKey::Type(ty_struct)]; + writeln!( + self.out, + "{struct_name} {}(in float arg) {{ + float whole; + float fract = modf(arg, whole); + {struct_name} result; + result.whole = whole; + result.fract = fract; + return result; +}}", + super::writer::MODF_FUNCTION, + )?; + 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..e47a57f98c 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, + crate::front::type_gen::FREXP_RESULT_F32_STRUCT_NAME, + super::writer::MODF_FUNCTION, + crate::front::type_gen::MODF_RESULT_F32_STRUCT_NAME, ]; // 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..b16320da65 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 FREXP_FUNCTION: &str = "naga_frexp"; +pub(crate) const MODF_FUNCTION: &str = "naga_modf"; + struct EpStructMember { name: String, ty: Handle, @@ -215,7 +218,10 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { // Write all structs for (handle, ty) in module.types.iter() { - if let TypeInner::Struct { ref members, span } = ty.inner { + if let TypeInner::Struct { + ref members, span, .. + } = ty.inner + { if module.types[members.last().unwrap().ty] .inner .is_dynamically_sized(&module.types) @@ -244,6 +250,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 +2673,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..22d51d7ea2 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, + crate::front::type_gen::FREXP_RESULT_F32_STRUCT_NAME, + super::writer::MODF_FUNCTION, + crate::front::type_gen::MODF_RESULT_F32_STRUCT_NAME, ]; diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index a8a103e6d0..ad4d104007 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 FREXP_FUNCTION: &str = "naga_frexp"; +pub(crate) const MODF_FUNCTION: &str = "naga_modf"; + /// 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,12 @@ 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).chain(arg1).chain(arg2).chain(arg3), + context, + )?; } else { write!(self.out, "{NAMESPACE}::{fun_name}")?; self.put_call_parameters( @@ -3236,6 +3245,31 @@ impl Writer { } } } + + if let Some(struct_ty) = module.special_types.frexp_result { + let struct_name = &self.names[&NameKey::Type(struct_ty)]; + writeln!( + self.out, + "struct {struct_name} {FREXP_FUNCTION}(float arg) {{ + int exp; + float fract = {NAMESPACE}::frexp(arg, exp); + return {struct_name}{{ fract, exp }}; +}};" + )?; + } + + if let Some(struct_ty) = module.special_types.modf_result { + let struct_name = &self.names[&NameKey::Type(struct_ty)]; + writeln!( + self.out, + "struct {struct_name} {MODF_FUNCTION}(float arg) {{ + float whole; + float fract = {NAMESPACE}::modf(arg, whole); + return {struct_name}{{ fract, whole }}; +}};" + )?; + } + 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/spv/writer.rs b/src/back/spv/writer.rs index e6db93602a..e0372f908b 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, @@ -1693,6 +1703,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 +1729,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 +1741,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/front/mod.rs b/src/front/mod.rs index 1f16ff6378..0900e5c86b 100644 --- a/src/front/mod.rs +++ b/src/front/mod.rs @@ -3,7 +3,7 @@ Frontend parsers that consume binary and text shaders and load them into [`Modul */ mod interpolator; -mod type_gen; +pub(crate) mod type_gen; #[cfg(feature = "glsl-in")] pub mod glsl; diff --git a/src/front/type_gen.rs b/src/front/type_gen.rs index 1ee454c448..170be9ddbc 100644 --- a/src/front/type_gen.rs +++ b/src/front/type_gen.rs @@ -4,6 +4,9 @@ Type generators. use crate::{arena::Handle, span::Span}; +pub(crate) const MODF_RESULT_F32_STRUCT_NAME: &str = "__modf_result_f32"; +pub(crate) const FREXP_RESULT_F32_STRUCT_NAME: &str = "__frexp_result_f32"; + impl crate::Module { pub fn generate_atomic_compare_exchange_result( &mut self, @@ -311,4 +314,101 @@ 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) { + if self.special_types.modf_result.is_some() { + return; + } + + let float_ty = self.types.insert( + crate::Type { + name: None, + inner: crate::TypeInner::Scalar { + kind: crate::ScalarKind::Float, + width: 4, + }, + }, + Span::UNDEFINED, + ); + + let handle = self.types.insert( + crate::Type { + name: Some("__modf_result_f32_".to_string()), + inner: crate::TypeInner::Struct { + members: vec![ + crate::StructMember { + name: Some("fract".to_string()), + ty: float_ty, + binding: None, + offset: 0, + }, + crate::StructMember { + name: Some("whole".to_string()), + ty: float_ty, + binding: None, + offset: 4, + }, + ], + span: 8, + }, + }, + Span::UNDEFINED, + ); + self.special_types.modf_result = Some(handle); + } + + /// Populate this module's [`crate::SpecialTypes::frexp_result`] type. + pub fn generate_frexp_result(&mut self) { + if self.special_types.frexp_result.is_some() { + return; + } + + let sint_ty = self.types.insert( + crate::Type { + name: None, + inner: crate::TypeInner::Scalar { + kind: crate::ScalarKind::Sint, + width: 4, + }, + }, + Span::UNDEFINED, + ); + + let float_ty = self.types.insert( + crate::Type { + name: None, + inner: crate::TypeInner::Scalar { + kind: crate::ScalarKind::Float, + width: 4, + }, + }, + Span::UNDEFINED, + ); + + let handle = self.types.insert( + crate::Type { + name: Some("__frexp_result_f32_".to_string()), + inner: crate::TypeInner::Struct { + members: vec![ + crate::StructMember { + name: Some("fract".to_string()), + ty: float_ty, + binding: None, + offset: 0, + }, + crate::StructMember { + name: Some("exp".to_string()), + ty: sint_ty, + binding: None, + offset: 4, + }, + ], + span: 8, + }, + }, + Span::UNDEFINED, + ); + self.special_types.frexp_result = Some(handle); + } } diff --git a/src/front/wgsl/lower/mod.rs b/src/front/wgsl/lower/mod.rs index f3a7b22dd3..776375d05d 100644 --- a/src/front/wgsl/lower/mod.rs +++ b/src/front/wgsl/lower/mod.rs @@ -1713,6 +1713,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { 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())) @@ -1731,6 +1732,12 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { args.finish()?; + if fun == crate::MathFunction::Frexp { + ctx.module.generate_frexp_result(); + } else if fun == crate::MathFunction::Modf { + ctx.module.generate_modf_result(); + }; + crate::Expression::Math { fun, arg, 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..c6e32e63bf 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -1961,6 +1961,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/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..905be978a0 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,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..95a0fe39ec 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -1015,33 +1015,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/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..c58dc73724 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 __modf_result_f32_ { + float fract_; + float whole; +}; +struct __frexp_result_f32_ { + float fract_; + int exp_; +}; + +__frexp_result_f32_ naga_frexp(float arg) { + int exp; + float fract = frexp(arg, exp); + return __frexp_result_f32_(fract, exp); +} + +__modf_result_f32_ naga_modf(float arg) { + float whole; + float fract = modf(arg, whole); + return __modf_result_f32_(fract, whole); +} 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))); + __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; + __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_; } diff --git a/tests/out/hlsl/math-functions.hlsl b/tests/out/hlsl/math-functions.hlsl index a1be810532..7ab294505d 100644 --- a/tests/out/hlsl/math-functions.hlsl +++ b/tests/out/hlsl/math-functions.hlsl @@ -1,3 +1,31 @@ +struct __modf_result_f32_ { + float fract; + float whole; +}; + +struct __frexp_result_f32_ { + float fract; + int exp_; +}; + +__frexp_result_f32_ naga_frexp(in float arg) { + float exp; + float fract = frexp(arg, exp); + __frexp_result_f32_ result; + result.exp = exp; + result.fract = fract; + return result; +} + +__modf_result_f32_ naga_modf(in float arg) { + float whole; + float fract = modf(arg, whole); + __modf_result_f32_ 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)); + __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; + __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_; } 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..090f59c999 100644 --- a/tests/out/msl/math-functions.msl +++ b/tests/out/msl/math-functions.msl @@ -4,6 +4,24 @@ using metal::uint; +struct __modf_result_f32_ { + float fract; + float whole; +}; +struct __frexp_result_f32_ { + float fract; + int exp; +}; +struct __frexp_result_f32_ naga_frexp(float arg) { + int exp; + float fract = metal::frexp(arg, exp); + return __frexp_result_f32_{ fract, exp }; +}; +struct __modf_result_f32_ naga_modf(float arg) { + float whole; + float fract = metal::modf(arg, whole); + return __modf_result_f32_{ fract, whole }; +}; fragment void main_( ) { @@ -38,4 +56,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)); + __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; + __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; } 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..4174b03518 100644 --- a/tests/out/wgsl/math-functions.wgsl +++ b/tests/out/wgsl/math-functions.wgsl @@ -1,3 +1,13 @@ +struct gen___modf_result_f32_ { + fract: f32, + whole: f32, +} + +struct gen___frexp_result_f32_ { + fract: f32, + exp: i32, +} + @fragment fn main() { let v = vec4(0.0); @@ -28,4 +38,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; }