From efa256c06224bad4ce9df82d7068ead9e1da1537 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 19 Jan 2024 14:07:22 -0500 Subject: [PATCH 1/7] refactor(const_eval): derive `Debug` for component-wise `enum`s --- naga/src/proc/constant_evaluator.rs | 1 + 1 file changed, 1 insertion(+) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index b5c821f412..6f70e13a2a 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -27,6 +27,7 @@ macro_rules! gen_component_wise_extractor { scalar_kinds: [$( $scalar_kind:ident ),* $(,)?], ) => { /// A subset of [`Literal`]s intended to be used for implementing numeric built-ins. + #[derive(Debug)] enum $target { $( #[doc = concat!( From cce2b5cf8552be5432b604da366e6e4c671011e8 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 19 Jan 2024 14:07:22 -0500 Subject: [PATCH 2/7] refactor(const_eval): derive `PartialEq` for testing in component-wise `enum`s --- naga/src/proc/constant_evaluator.rs | 1 + 1 file changed, 1 insertion(+) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 6f70e13a2a..cd7bb9f34b 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -28,6 +28,7 @@ macro_rules! gen_component_wise_extractor { ) => { /// A subset of [`Literal`]s intended to be used for implementing numeric built-ins. #[derive(Debug)] + #[cfg_attr(test, derive(PartialEq))] enum $target { $( #[doc = concat!( From f46fc01abb6dcc280380520a85815dcfe69f9991 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 19 Jan 2024 14:04:57 -0500 Subject: [PATCH 3/7] refactor(naga): rename `MathFunction::FindMsb` to `FirstLeadingBit` --- naga/src/back/glsl/mod.rs | 8 +++++--- naga/src/back/hlsl/writer.rs | 2 +- naga/src/back/msl/writer.rs | 6 +++--- naga/src/back/spv/block.rs | 4 ++-- naga/src/back/wgsl/writer.rs | 2 +- naga/src/front/glsl/builtins.rs | 8 +++++--- naga/src/front/spv/mod.rs | 2 +- naga/src/front/wgsl/parse/conv.rs | 2 +- naga/src/lib.rs | 2 +- naga/src/proc/mod.rs | 2 +- naga/src/proc/typifier.rs | 2 +- naga/src/valid/expression.rs | 2 +- 12 files changed, 23 insertions(+), 19 deletions(-) diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index 7ad1f3c597..9ea6eed91a 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -3648,7 +3648,7 @@ impl<'a, W: Write> Writer<'a, W> { return Ok(()); } Mf::FindLsb => "findLSB", - Mf::FindMsb => "findMSB", + Mf::FirstLeadingBit => "findMSB", // data packing Mf::Pack4x8snorm => "packSnorm4x8", Mf::Pack4x8unorm => "packUnorm4x8", @@ -3722,8 +3722,10 @@ impl<'a, W: Write> Writer<'a, W> { // Some GLSL functions always return signed integers (like findMSB), // so they need to be cast to uint if the argument is also an uint. - let ret_might_need_int_to_uint = - matches!(fun, Mf::FindLsb | Mf::FindMsb | Mf::CountOneBits | Mf::Abs); + let ret_might_need_int_to_uint = matches!( + fun, + Mf::FindLsb | Mf::FirstLeadingBit | Mf::CountOneBits | Mf::Abs + ); // Some GLSL functions only accept signed integers (like abs), // so they need their argument cast from uint to int. diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index 982bf0cfea..7965e6492b 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/naga/src/back/hlsl/writer.rs @@ -3064,7 +3064,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { Mf::CountOneBits => Function::MissingIntOverload("countbits"), Mf::ReverseBits => Function::MissingIntOverload("reversebits"), Mf::FindLsb => Function::MissingIntReturnType("firstbitlow"), - Mf::FindMsb => Function::MissingIntReturnType("firstbithigh"), + Mf::FirstLeadingBit => Function::MissingIntReturnType("firstbithigh"), Mf::ExtractBits => Function::Regular(EXTRACT_BITS_FUNCTION), Mf::InsertBits => Function::Regular(INSERT_BITS_FUNCTION), // Data Packing diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index 7ec22009bd..fccc92a1db 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -1876,7 +1876,7 @@ impl Writer { Mf::ExtractBits => "", Mf::InsertBits => "", Mf::FindLsb => "", - Mf::FindMsb => "", + Mf::FirstLeadingBit => "", // data packing Mf::Pack4x8snorm => "pack_float_to_snorm4x8", Mf::Pack4x8unorm => "pack_float_to_unorm4x8", @@ -1928,7 +1928,7 @@ impl Writer { self.put_expression(arg, context, true)?; write!(self.out, ") + 1) % {constant}) - 1)")?; } - Mf::FindMsb => { + Mf::FirstLeadingBit => { let inner = context.resolve_type(arg); let scalar = inner.scalar().unwrap(); let constant = scalar.width * 8 - 1; @@ -2702,7 +2702,7 @@ impl Writer { } } } - crate::MathFunction::FindMsb + crate::MathFunction::FirstLeadingBit | crate::MathFunction::Pack4xI8 | crate::MathFunction::Pack4xU8 | crate::MathFunction::Unpack4xI8 diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index 33f892aa45..932e27cceb 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -1184,12 +1184,12 @@ impl<'w> BlockContext<'w> { )) } Mf::FindLsb => MathOp::Ext(spirv::GLOp::FindILsb), - Mf::FindMsb => { + Mf::FirstLeadingBit => { if arg_ty.scalar_width() == Some(4) { let thing = match arg_scalar_kind { Some(crate::ScalarKind::Uint) => spirv::GLOp::FindUMsb, Some(crate::ScalarKind::Sint) => spirv::GLOp::FindSMsb, - other => unimplemented!("Unexpected findMSB({:?})", other), + other => unimplemented!("Unexpected firstLeadingBit({:?})", other), }; MathOp::Ext(thing) } else { diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs index 8cd37830ec..1b3597bcba 100644 --- a/naga/src/back/wgsl/writer.rs +++ b/naga/src/back/wgsl/writer.rs @@ -1711,7 +1711,7 @@ impl Writer { Mf::ExtractBits => Function::Regular("extractBits"), Mf::InsertBits => Function::Regular("insertBits"), Mf::FindLsb => Function::Regular("firstTrailingBit"), - Mf::FindMsb => Function::Regular("firstLeadingBit"), + Mf::FirstLeadingBit => Function::Regular("firstLeadingBit"), // data packing Mf::Pack4x8snorm => Function::Regular("pack4x8snorm"), Mf::Pack4x8unorm => Function::Regular("pack4x8unorm"), diff --git a/naga/src/front/glsl/builtins.rs b/naga/src/front/glsl/builtins.rs index cbb9b99387..b0e921b79c 100644 --- a/naga/src/front/glsl/builtins.rs +++ b/naga/src/front/glsl/builtins.rs @@ -647,7 +647,7 @@ fn inject_standard_builtins( "bitfieldExtract" => MathFunction::ExtractBits, "bitfieldInsert" => MathFunction::InsertBits, "findLSB" => MathFunction::FindLsb, - "findMSB" => MathFunction::FindMsb, + "findMSB" => MathFunction::FirstLeadingBit, _ => unreachable!(), }; @@ -696,7 +696,9 @@ fn inject_standard_builtins( let mc = if scalar.kind == Sk::Uint { match mc { MacroCall::MathFunction(MathFunction::FindLsb) => MacroCall::FindLsbUint, - MacroCall::MathFunction(MathFunction::FindMsb) => MacroCall::FindMsbUint, + MacroCall::MathFunction(MathFunction::FirstLeadingBit) => { + MacroCall::FindMsbUint + } mc => mc, } } else { @@ -1788,7 +1790,7 @@ impl MacroCall { mc @ (MacroCall::FindLsbUint | MacroCall::FindMsbUint) => { let fun = match mc { MacroCall::FindLsbUint => MathFunction::FindLsb, - MacroCall::FindMsbUint => MathFunction::FindMsb, + MacroCall::FindMsbUint => MathFunction::FirstLeadingBit, _ => unreachable!(), }; let res = ctx.add_expression( diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs index d154712b20..ac048203e4 100644 --- a/naga/src/front/spv/mod.rs +++ b/naga/src/front/spv/mod.rs @@ -3027,7 +3027,7 @@ impl> Frontend { Glo::UnpackUnorm2x16 => Mf::Unpack2x16unorm, Glo::UnpackSnorm2x16 => Mf::Unpack2x16snorm, Glo::FindILsb => Mf::FindLsb, - Glo::FindUMsb | Glo::FindSMsb => Mf::FindMsb, + Glo::FindUMsb | Glo::FindSMsb => Mf::FirstLeadingBit, // TODO: https://github.com/gfx-rs/naga/issues/2526 Glo::Modf | Glo::Frexp => return Err(Error::UnsupportedExtInst(inst_id)), Glo::IMix diff --git a/naga/src/front/wgsl/parse/conv.rs b/naga/src/front/wgsl/parse/conv.rs index 49b15dfa83..2cb676a80d 100644 --- a/naga/src/front/wgsl/parse/conv.rs +++ b/naga/src/front/wgsl/parse/conv.rs @@ -236,7 +236,7 @@ pub fn map_standard_fun(word: &str) -> Option { "extractBits" => Mf::ExtractBits, "insertBits" => Mf::InsertBits, "firstTrailingBit" => Mf::FindLsb, - "firstLeadingBit" => Mf::FindMsb, + "firstLeadingBit" => Mf::FirstLeadingBit, // data packing "pack4x8snorm" => Mf::Pack4x8snorm, "pack4x8unorm" => Mf::Pack4x8unorm, diff --git a/naga/src/lib.rs b/naga/src/lib.rs index 8ed7527922..8a52df81bb 100644 --- a/naga/src/lib.rs +++ b/naga/src/lib.rs @@ -1199,7 +1199,7 @@ pub enum MathFunction { ExtractBits, InsertBits, FindLsb, - FindMsb, + FirstLeadingBit, // data packing Pack4x8snorm, Pack4x8unorm, diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs index 86d2b11f25..6b514af18b 100644 --- a/naga/src/proc/mod.rs +++ b/naga/src/proc/mod.rs @@ -485,7 +485,7 @@ impl super::MathFunction { Self::ExtractBits => 3, Self::InsertBits => 4, Self::FindLsb => 1, - Self::FindMsb => 1, + Self::FirstLeadingBit => 1, // data packing Self::Pack4x8snorm => 1, Self::Pack4x8unorm => 1, diff --git a/naga/src/proc/typifier.rs b/naga/src/proc/typifier.rs index 0a02900c4a..23295cc0e2 100644 --- a/naga/src/proc/typifier.rs +++ b/naga/src/proc/typifier.rs @@ -789,7 +789,7 @@ impl<'a> ResolveContext<'a> { Mf::ExtractBits | Mf::InsertBits | Mf::FindLsb | - Mf::FindMsb => match *res_arg.inner_with(types) { + Mf::FirstLeadingBit => match *res_arg.inner_with(types) { Ti::Scalar(scalar @ crate::Scalar { kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint, .. diff --git a/naga/src/valid/expression.rs b/naga/src/valid/expression.rs index 89bceae061..bd90c8ad2d 100644 --- a/naga/src/valid/expression.rs +++ b/naga/src/valid/expression.rs @@ -1350,7 +1350,7 @@ impl super::Validator { | Mf::CountTrailingZeros | Mf::CountOneBits | Mf::ReverseBits - | Mf::FindMsb + | Mf::FirstLeadingBit | Mf::FindLsb => { if arg1_ty.is_some() || arg2_ty.is_some() || arg3_ty.is_some() { return Err(ExpressionError::WrongArgumentCount(fun)); From c7ab27c72e6f7cf92fce5ecf8f13715054355fe0 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 19 Jan 2024 14:07:22 -0500 Subject: [PATCH 4/7] feat(const_eval): impl. `firstLeadingBit` --- naga/src/proc/constant_evaluator.rs | 91 ++++++++++++ .../glsl/math-functions.main.Fragment.glsl | 6 +- naga/tests/out/hlsl/math-functions.hlsl | 6 +- naga/tests/out/msl/math-functions.msl | 8 +- naga/tests/out/spv/math-functions.spvasm | 136 +++++++++--------- naga/tests/out/wgsl/math-functions.wgsl | 6 +- 6 files changed, 166 insertions(+), 87 deletions(-) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index cd7bb9f34b..344e2ddba2 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -1233,6 +1233,9 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::ReverseBits => { component_wise_concrete_int!(self, span, [arg], |e| { Ok([e.reverse_bits()]) }) } + crate::MathFunction::FirstLeadingBit => { + component_wise_concrete_int(self, span, [arg], |ci| Ok(first_leading_bit(ci))) + } fun => Err(ConstantEvaluatorError::NotImplemented(format!( "{fun:?} built-in function" @@ -2098,6 +2101,94 @@ impl<'a> ConstantEvaluator<'a> { } } +fn first_leading_bit(concrete_int: ConcreteInt<1>) -> ConcreteInt<1> { + // NOTE: Bit indices for this built-in start at 0 at the "right" (or LSB). For example, 1 means + // the least significant bit is set. Therefore, an input of 1 would return a right-to-left bit + // index of 0. + let rtl_to_ltr_bit_idx = |e: u32| -> u32 { + match e { + idx @ 0..=31 => 31 - idx, + 32 => u32::MAX, + _ => unreachable!(), + } + }; + match concrete_int { + ConcreteInt::I32([e]) => ConcreteInt::I32([{ + let rtl_bit_index = if e.is_negative() { + e.leading_ones() + } else { + e.leading_zeros() + }; + rtl_to_ltr_bit_idx(rtl_bit_index) as i32 + }]), + ConcreteInt::U32([e]) => ConcreteInt::U32([rtl_to_ltr_bit_idx(e.leading_zeros())]), + } +} + +#[test] +fn first_leading_bit_smoke() { + assert_eq!( + first_leading_bit(ConcreteInt::I32([-1])), + ConcreteInt::I32([-1]) + ); + assert_eq!( + first_leading_bit(ConcreteInt::I32([0])), + ConcreteInt::I32([-1]) + ); + assert_eq!( + first_leading_bit(ConcreteInt::I32([1])), + ConcreteInt::I32([0]) + ); + assert_eq!( + first_leading_bit(ConcreteInt::I32([-2])), + ConcreteInt::I32([0]) + ); + assert_eq!( + first_leading_bit(ConcreteInt::I32([1234 + 4567])), + ConcreteInt::I32([12]) + ); + assert_eq!( + first_leading_bit(ConcreteInt::I32([i32::MAX])), + ConcreteInt::I32([30]) + ); + assert_eq!( + first_leading_bit(ConcreteInt::I32([i32::MIN])), + ConcreteInt::I32([30]) + ); + // NOTE: Ignore the sign bit, which is a separate (above) case. + for idx in 0..(32 - 1) { + assert_eq!( + first_leading_bit(ConcreteInt::I32([1 << idx])), + ConcreteInt::I32([idx]) + ); + } + for idx in 1..(32 - 1) { + assert_eq!( + first_leading_bit(ConcreteInt::I32([-(1 << idx)])), + ConcreteInt::I32([idx - 1]) + ); + } + + assert_eq!( + first_leading_bit(ConcreteInt::U32([0])), + ConcreteInt::U32([u32::MAX]) + ); + assert_eq!( + first_leading_bit(ConcreteInt::U32([1])), + ConcreteInt::U32([0]) + ); + assert_eq!( + first_leading_bit(ConcreteInt::U32([u32::MAX])), + ConcreteInt::U32([31]) + ); + for idx in 0..32 { + assert_eq!( + first_leading_bit(ConcreteInt::U32([1 << idx])), + ConcreteInt::U32([idx]) + ) + } +} + /// Trait for conversions of abstract values to concrete types. trait TryFromAbstract: Sized { /// Convert an abstract literal `value` to `Self`. diff --git a/naga/tests/out/glsl/math-functions.main.Fragment.glsl b/naga/tests/out/glsl/math-functions.main.Fragment.glsl index 7f91571dcc..c10dddff0f 100644 --- a/naga/tests/out/glsl/math-functions.main.Fragment.glsl +++ b/naga/tests/out/glsl/math-functions.main.Fragment.glsl @@ -65,10 +65,8 @@ void main() { ivec4 sign_b = ivec4(-1, -1, -1, -1); vec4 sign_d = vec4(-1.0, -1.0, -1.0, -1.0); int const_dot = ( + ivec2(0).x * ivec2(0).x + ivec2(0).y * ivec2(0).y); - uint first_leading_bit_abs = uint(findMSB(0u)); - int flb_a = findMSB(-1); - ivec2 flb_b = findMSB(ivec2(-1)); - uvec2 flb_c = uvec2(findMSB(uvec2(1u))); + ivec2 flb_b = ivec2(-1, -1); + uvec2 flb_c = uvec2(0u, 0u); int ftb_a = findLSB(-1); uint ftb_b = uint(findLSB(1u)); ivec2 ftb_c = findLSB(ivec2(-1)); diff --git a/naga/tests/out/hlsl/math-functions.hlsl b/naga/tests/out/hlsl/math-functions.hlsl index c1a771c25d..14d1e9e188 100644 --- a/naga/tests/out/hlsl/math-functions.hlsl +++ b/naga/tests/out/hlsl/math-functions.hlsl @@ -79,10 +79,8 @@ void main() int4 sign_b = int4(-1, -1, -1, -1); float4 sign_d = float4(-1.0, -1.0, -1.0, -1.0); int const_dot = dot(ZeroValueint2(), ZeroValueint2()); - uint first_leading_bit_abs = firstbithigh(0u); - int flb_a = asint(firstbithigh(-1)); - int2 flb_b = asint(firstbithigh((-1).xx)); - uint2 flb_c = firstbithigh((1u).xx); + int2 flb_b = int2(-1, -1); + uint2 flb_c = uint2(0u, 0u); int ftb_a = asint(firstbitlow(-1)); uint ftb_b = firstbitlow(1u); int2 ftb_c = asint(firstbitlow((-1).xx)); diff --git a/naga/tests/out/msl/math-functions.msl b/naga/tests/out/msl/math-functions.msl index 0e6a5b24dc..271472978a 100644 --- a/naga/tests/out/msl/math-functions.msl +++ b/naga/tests/out/msl/math-functions.msl @@ -67,12 +67,8 @@ fragment void main_( metal::int4 sign_b = metal::int4(-1, -1, -1, -1); metal::float4 sign_d = metal::float4(-1.0, -1.0, -1.0, -1.0); int const_dot = ( + metal::int2 {}.x * metal::int2 {}.x + metal::int2 {}.y * metal::int2 {}.y); - uint first_leading_bit_abs = metal::select(31 - metal::clz(0u), uint(-1), 0u == 0 || 0u == -1); - int flb_a = metal::select(31 - metal::clz(metal::select(-1, ~-1, -1 < 0)), int(-1), -1 == 0 || -1 == -1); - metal::int2 _e29 = metal::int2(-1); - metal::int2 flb_b = metal::select(31 - metal::clz(metal::select(_e29, ~_e29, _e29 < 0)), int2(-1), _e29 == 0 || _e29 == -1); - metal::uint2 _e32 = metal::uint2(1u); - metal::uint2 flb_c = metal::select(31 - metal::clz(_e32), uint2(-1), _e32 == 0 || _e32 == -1); + metal::int2 flb_b = metal::int2(-1, -1); + metal::uint2 flb_c = metal::uint2(0u, 0u); int ftb_a = (((metal::ctz(-1) + 1) % 33) - 1); uint ftb_b = (((metal::ctz(1u) + 1) % 33) - 1); metal::int2 ftb_c = (((metal::ctz(metal::int2(-1)) + 1) % 33) - 1); diff --git a/naga/tests/out/spv/math-functions.spvasm b/naga/tests/out/spv/math-functions.spvasm index 6e07c6d7a6..2207934cc9 100644 --- a/naga/tests/out/spv/math-functions.spvasm +++ b/naga/tests/out/spv/math-functions.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 96 +; Bound: 94 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 @@ -40,77 +40,75 @@ OpMemberDecorate %15 1 Offset 16 %24 = OpConstant %4 -1.0 %25 = OpConstantComposite %3 %24 %24 %24 %24 %26 = OpConstantNull %7 -%27 = OpConstant %9 0 +%27 = OpConstant %9 4294967295 %28 = OpConstantComposite %7 %22 %22 -%29 = OpConstant %9 1 +%29 = OpConstant %9 0 %30 = OpConstantComposite %8 %29 %29 -%31 = OpConstant %9 32 -%32 = OpConstant %6 32 -%33 = OpConstant %6 0 -%34 = OpConstantComposite %8 %31 %31 -%35 = OpConstantComposite %7 %32 %32 -%36 = OpConstantComposite %8 %27 %27 -%37 = OpConstantComposite %7 %33 %33 -%38 = OpConstant %9 31 -%39 = OpConstantComposite %8 %38 %38 -%40 = OpConstant %6 2 -%41 = OpConstant %4 2.0 -%42 = OpConstantComposite %10 %19 %41 -%43 = OpConstant %6 3 -%44 = OpConstant %6 4 -%45 = OpConstantComposite %7 %43 %44 -%46 = OpConstant %4 1.5 -%47 = OpConstantComposite %10 %46 %46 -%48 = OpConstantComposite %3 %46 %46 %46 %46 -%55 = OpConstantComposite %3 %19 %19 %19 %19 -%58 = OpConstantNull %6 +%31 = OpConstant %9 1 +%32 = OpConstantComposite %7 %22 %22 +%33 = OpConstantComposite %8 %31 %31 +%34 = OpConstant %9 32 +%35 = OpConstant %6 32 +%36 = OpConstant %6 0 +%37 = OpConstantComposite %8 %34 %34 +%38 = OpConstantComposite %7 %35 %35 +%39 = OpConstantComposite %7 %36 %36 +%40 = OpConstant %9 31 +%41 = OpConstantComposite %8 %40 %40 +%42 = OpConstant %6 2 +%43 = OpConstant %4 2.0 +%44 = OpConstantComposite %10 %19 %43 +%45 = OpConstant %6 3 +%46 = OpConstant %6 4 +%47 = OpConstantComposite %7 %45 %46 +%48 = OpConstant %4 1.5 +%49 = OpConstantComposite %10 %48 %48 +%50 = OpConstantComposite %3 %48 %48 %48 %48 +%57 = OpConstantComposite %3 %19 %19 %19 %19 +%60 = OpConstantNull %6 %17 = OpFunction %2 None %18 %16 = OpLabel -OpBranch %49 -%49 = OpLabel -%50 = OpExtInst %4 %1 Degrees %19 -%51 = OpExtInst %4 %1 Radians %19 -%52 = OpExtInst %3 %1 Degrees %21 -%53 = OpExtInst %3 %1 Radians %21 -%54 = OpExtInst %3 %1 FClamp %21 %21 %55 -%56 = OpExtInst %3 %1 Refract %21 %21 %19 -%59 = OpCompositeExtract %6 %26 0 -%60 = OpCompositeExtract %6 %26 0 -%61 = OpIMul %6 %59 %60 -%62 = OpIAdd %6 %58 %61 -%63 = OpCompositeExtract %6 %26 1 -%64 = OpCompositeExtract %6 %26 1 -%65 = OpIMul %6 %63 %64 -%57 = OpIAdd %6 %62 %65 -%66 = OpExtInst %9 %1 FindUMsb %27 -%67 = OpExtInst %6 %1 FindSMsb %22 -%68 = OpExtInst %7 %1 FindSMsb %28 -%69 = OpExtInst %8 %1 FindUMsb %30 -%70 = OpExtInst %6 %1 FindILsb %22 -%71 = OpExtInst %9 %1 FindILsb %29 -%72 = OpExtInst %7 %1 FindILsb %28 -%73 = OpExtInst %8 %1 FindILsb %30 -%74 = OpExtInst %4 %1 Ldexp %19 %40 -%75 = OpExtInst %10 %1 Ldexp %42 %45 -%76 = OpExtInst %11 %1 ModfStruct %46 -%77 = OpExtInst %11 %1 ModfStruct %46 -%78 = OpCompositeExtract %4 %77 0 -%79 = OpExtInst %11 %1 ModfStruct %46 -%80 = OpCompositeExtract %4 %79 1 -%81 = OpExtInst %12 %1 ModfStruct %47 -%82 = OpExtInst %13 %1 ModfStruct %48 -%83 = OpCompositeExtract %3 %82 1 -%84 = OpCompositeExtract %4 %83 0 -%85 = OpExtInst %12 %1 ModfStruct %47 -%86 = OpCompositeExtract %10 %85 0 -%87 = OpCompositeExtract %4 %86 1 -%88 = OpExtInst %14 %1 FrexpStruct %46 -%89 = OpExtInst %14 %1 FrexpStruct %46 -%90 = OpCompositeExtract %4 %89 0 -%91 = OpExtInst %14 %1 FrexpStruct %46 -%92 = OpCompositeExtract %6 %91 1 -%93 = OpExtInst %15 %1 FrexpStruct %48 -%94 = OpCompositeExtract %5 %93 1 -%95 = OpCompositeExtract %6 %94 0 +OpBranch %51 +%51 = OpLabel +%52 = OpExtInst %4 %1 Degrees %19 +%53 = OpExtInst %4 %1 Radians %19 +%54 = OpExtInst %3 %1 Degrees %21 +%55 = OpExtInst %3 %1 Radians %21 +%56 = OpExtInst %3 %1 FClamp %21 %21 %57 +%58 = OpExtInst %3 %1 Refract %21 %21 %19 +%61 = OpCompositeExtract %6 %26 0 +%62 = OpCompositeExtract %6 %26 0 +%63 = OpIMul %6 %61 %62 +%64 = OpIAdd %6 %60 %63 +%65 = OpCompositeExtract %6 %26 1 +%66 = OpCompositeExtract %6 %26 1 +%67 = OpIMul %6 %65 %66 +%59 = OpIAdd %6 %64 %67 +%68 = OpExtInst %6 %1 FindILsb %22 +%69 = OpExtInst %9 %1 FindILsb %31 +%70 = OpExtInst %7 %1 FindILsb %32 +%71 = OpExtInst %8 %1 FindILsb %33 +%72 = OpExtInst %4 %1 Ldexp %19 %42 +%73 = OpExtInst %10 %1 Ldexp %44 %47 +%74 = OpExtInst %11 %1 ModfStruct %48 +%75 = OpExtInst %11 %1 ModfStruct %48 +%76 = OpCompositeExtract %4 %75 0 +%77 = OpExtInst %11 %1 ModfStruct %48 +%78 = OpCompositeExtract %4 %77 1 +%79 = OpExtInst %12 %1 ModfStruct %49 +%80 = OpExtInst %13 %1 ModfStruct %50 +%81 = OpCompositeExtract %3 %80 1 +%82 = OpCompositeExtract %4 %81 0 +%83 = OpExtInst %12 %1 ModfStruct %49 +%84 = OpCompositeExtract %10 %83 0 +%85 = OpCompositeExtract %4 %84 1 +%86 = OpExtInst %14 %1 FrexpStruct %48 +%87 = OpExtInst %14 %1 FrexpStruct %48 +%88 = OpCompositeExtract %4 %87 0 +%89 = OpExtInst %14 %1 FrexpStruct %48 +%90 = OpCompositeExtract %6 %89 1 +%91 = OpExtInst %15 %1 FrexpStruct %50 +%92 = OpCompositeExtract %5 %91 1 +%93 = OpCompositeExtract %6 %92 0 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/math-functions.wgsl b/naga/tests/out/wgsl/math-functions.wgsl index 228248b3ce..c97ce6a4a2 100644 --- a/naga/tests/out/wgsl/math-functions.wgsl +++ b/naga/tests/out/wgsl/math-functions.wgsl @@ -10,10 +10,8 @@ fn main() { let sign_b = vec4(-1i, -1i, -1i, -1i); let sign_d = vec4(-1f, -1f, -1f, -1f); let const_dot = dot(vec2(), vec2()); - let first_leading_bit_abs = firstLeadingBit(0u); - let flb_a = firstLeadingBit(-1i); - let flb_b = firstLeadingBit(vec2(-1i)); - let flb_c = firstLeadingBit(vec2(1u)); + let flb_b = vec2(-1i, -1i); + let flb_c = vec2(0u, 0u); let ftb_a = firstTrailingBit(-1i); let ftb_b = firstTrailingBit(1u); let ftb_c = firstTrailingBit(vec2(-1i)); From c6dc57e35440e3dac0f6281eaa7bdcefc647b971 Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 19 Jan 2024 14:15:49 -0500 Subject: [PATCH 5/7] refactor(naga): rename `MathFunction::FindLsb` to `FirstTrailingBit` --- naga/src/back/glsl/mod.rs | 4 ++-- naga/src/back/hlsl/writer.rs | 2 +- naga/src/back/msl/writer.rs | 4 ++-- naga/src/back/spv/block.rs | 2 +- naga/src/back/wgsl/writer.rs | 2 +- naga/src/front/glsl/builtins.rs | 8 +++++--- naga/src/front/spv/mod.rs | 2 +- naga/src/front/wgsl/parse/conv.rs | 2 +- naga/src/lib.rs | 2 +- naga/src/proc/mod.rs | 2 +- naga/src/proc/typifier.rs | 2 +- naga/src/valid/expression.rs | 2 +- 12 files changed, 18 insertions(+), 16 deletions(-) diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index 9ea6eed91a..fe70480544 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -3647,7 +3647,7 @@ impl<'a, W: Write> Writer<'a, W> { return Ok(()); } - Mf::FindLsb => "findLSB", + Mf::FirstTrailingBit => "findLSB", Mf::FirstLeadingBit => "findMSB", // data packing Mf::Pack4x8snorm => "packSnorm4x8", @@ -3724,7 +3724,7 @@ impl<'a, W: Write> Writer<'a, W> { // so they need to be cast to uint if the argument is also an uint. let ret_might_need_int_to_uint = matches!( fun, - Mf::FindLsb | Mf::FirstLeadingBit | Mf::CountOneBits | Mf::Abs + Mf::FirstTrailingBit | Mf::FirstLeadingBit | Mf::CountOneBits | Mf::Abs ); // Some GLSL functions only accept signed integers (like abs), diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index 7965e6492b..85d943e850 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/naga/src/back/hlsl/writer.rs @@ -3063,7 +3063,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { Mf::CountLeadingZeros => Function::CountLeadingZeros, Mf::CountOneBits => Function::MissingIntOverload("countbits"), Mf::ReverseBits => Function::MissingIntOverload("reversebits"), - Mf::FindLsb => Function::MissingIntReturnType("firstbitlow"), + Mf::FirstTrailingBit => Function::MissingIntReturnType("firstbitlow"), Mf::FirstLeadingBit => Function::MissingIntReturnType("firstbithigh"), Mf::ExtractBits => Function::Regular(EXTRACT_BITS_FUNCTION), Mf::InsertBits => Function::Regular(INSERT_BITS_FUNCTION), diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index fccc92a1db..b112bb369a 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -1875,7 +1875,7 @@ impl Writer { Mf::ReverseBits => "reverse_bits", Mf::ExtractBits => "", Mf::InsertBits => "", - Mf::FindLsb => "", + Mf::FirstTrailingBit => "", Mf::FirstLeadingBit => "", // data packing Mf::Pack4x8snorm => "pack_float_to_snorm4x8", @@ -1920,7 +1920,7 @@ impl Writer { self.put_expression(arg1.unwrap(), context, false)?; write!(self.out, ")")?; } - Mf::FindLsb => { + Mf::FirstTrailingBit => { let scalar = context.resolve_type(arg).scalar().unwrap(); let constant = scalar.width * 8 + 1; diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index 932e27cceb..9fb9485860 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -1183,7 +1183,7 @@ impl<'w> BlockContext<'w> { count_id, )) } - Mf::FindLsb => MathOp::Ext(spirv::GLOp::FindILsb), + Mf::FirstTrailingBit => MathOp::Ext(spirv::GLOp::FindILsb), Mf::FirstLeadingBit => { if arg_ty.scalar_width() == Some(4) { let thing = match arg_scalar_kind { diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs index 1b3597bcba..6a069113eb 100644 --- a/naga/src/back/wgsl/writer.rs +++ b/naga/src/back/wgsl/writer.rs @@ -1710,7 +1710,7 @@ impl Writer { Mf::ReverseBits => Function::Regular("reverseBits"), Mf::ExtractBits => Function::Regular("extractBits"), Mf::InsertBits => Function::Regular("insertBits"), - Mf::FindLsb => Function::Regular("firstTrailingBit"), + Mf::FirstTrailingBit => Function::Regular("firstTrailingBit"), Mf::FirstLeadingBit => Function::Regular("firstLeadingBit"), // data packing Mf::Pack4x8snorm => Function::Regular("pack4x8snorm"), diff --git a/naga/src/front/glsl/builtins.rs b/naga/src/front/glsl/builtins.rs index b0e921b79c..f76ce7754a 100644 --- a/naga/src/front/glsl/builtins.rs +++ b/naga/src/front/glsl/builtins.rs @@ -646,7 +646,7 @@ fn inject_standard_builtins( "bitfieldReverse" => MathFunction::ReverseBits, "bitfieldExtract" => MathFunction::ExtractBits, "bitfieldInsert" => MathFunction::InsertBits, - "findLSB" => MathFunction::FindLsb, + "findLSB" => MathFunction::FirstTrailingBit, "findMSB" => MathFunction::FirstLeadingBit, _ => unreachable!(), }; @@ -695,7 +695,9 @@ fn inject_standard_builtins( // we need to cast the return type of findLsb / findMsb let mc = if scalar.kind == Sk::Uint { match mc { - MacroCall::MathFunction(MathFunction::FindLsb) => MacroCall::FindLsbUint, + MacroCall::MathFunction(MathFunction::FirstTrailingBit) => { + MacroCall::FindLsbUint + } MacroCall::MathFunction(MathFunction::FirstLeadingBit) => { MacroCall::FindMsbUint } @@ -1789,7 +1791,7 @@ impl MacroCall { )?, mc @ (MacroCall::FindLsbUint | MacroCall::FindMsbUint) => { let fun = match mc { - MacroCall::FindLsbUint => MathFunction::FindLsb, + MacroCall::FindLsbUint => MathFunction::FirstTrailingBit, MacroCall::FindMsbUint => MathFunction::FirstLeadingBit, _ => unreachable!(), }; diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs index ac048203e4..809aff7674 100644 --- a/naga/src/front/spv/mod.rs +++ b/naga/src/front/spv/mod.rs @@ -3026,7 +3026,7 @@ impl> Frontend { Glo::UnpackHalf2x16 => Mf::Unpack2x16float, Glo::UnpackUnorm2x16 => Mf::Unpack2x16unorm, Glo::UnpackSnorm2x16 => Mf::Unpack2x16snorm, - Glo::FindILsb => Mf::FindLsb, + Glo::FindILsb => Mf::FirstTrailingBit, Glo::FindUMsb | Glo::FindSMsb => Mf::FirstLeadingBit, // TODO: https://github.com/gfx-rs/naga/issues/2526 Glo::Modf | Glo::Frexp => return Err(Error::UnsupportedExtInst(inst_id)), diff --git a/naga/src/front/wgsl/parse/conv.rs b/naga/src/front/wgsl/parse/conv.rs index 2cb676a80d..80f05db59a 100644 --- a/naga/src/front/wgsl/parse/conv.rs +++ b/naga/src/front/wgsl/parse/conv.rs @@ -235,7 +235,7 @@ pub fn map_standard_fun(word: &str) -> Option { "reverseBits" => Mf::ReverseBits, "extractBits" => Mf::ExtractBits, "insertBits" => Mf::InsertBits, - "firstTrailingBit" => Mf::FindLsb, + "firstTrailingBit" => Mf::FirstTrailingBit, "firstLeadingBit" => Mf::FirstLeadingBit, // data packing "pack4x8snorm" => Mf::Pack4x8snorm, diff --git a/naga/src/lib.rs b/naga/src/lib.rs index 8a52df81bb..94edec9159 100644 --- a/naga/src/lib.rs +++ b/naga/src/lib.rs @@ -1198,7 +1198,7 @@ pub enum MathFunction { ReverseBits, ExtractBits, InsertBits, - FindLsb, + FirstTrailingBit, FirstLeadingBit, // data packing Pack4x8snorm, diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs index 6b514af18b..41273c5c72 100644 --- a/naga/src/proc/mod.rs +++ b/naga/src/proc/mod.rs @@ -484,7 +484,7 @@ impl super::MathFunction { Self::ReverseBits => 1, Self::ExtractBits => 3, Self::InsertBits => 4, - Self::FindLsb => 1, + Self::FirstTrailingBit => 1, Self::FirstLeadingBit => 1, // data packing Self::Pack4x8snorm => 1, diff --git a/naga/src/proc/typifier.rs b/naga/src/proc/typifier.rs index 23295cc0e2..d8af0cd236 100644 --- a/naga/src/proc/typifier.rs +++ b/naga/src/proc/typifier.rs @@ -788,7 +788,7 @@ impl<'a> ResolveContext<'a> { Mf::ReverseBits | Mf::ExtractBits | Mf::InsertBits | - Mf::FindLsb | + Mf::FirstTrailingBit | Mf::FirstLeadingBit => match *res_arg.inner_with(types) { Ti::Scalar(scalar @ crate::Scalar { kind: crate::ScalarKind::Sint | crate::ScalarKind::Uint, diff --git a/naga/src/valid/expression.rs b/naga/src/valid/expression.rs index bd90c8ad2d..116560bb61 100644 --- a/naga/src/valid/expression.rs +++ b/naga/src/valid/expression.rs @@ -1351,7 +1351,7 @@ impl super::Validator { | Mf::CountOneBits | Mf::ReverseBits | Mf::FirstLeadingBit - | Mf::FindLsb => { + | Mf::FirstTrailingBit => { if arg1_ty.is_some() || arg2_ty.is_some() || arg3_ty.is_some() { return Err(ExpressionError::WrongArgumentCount(fun)); } From d122588944ef1d785928c327123aafc3a9dd066a Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Fri, 26 Jan 2024 13:06:20 -0500 Subject: [PATCH 6/7] feat(const_eval): impl. `firstTrailingBit` --- naga/src/proc/constant_evaluator.rs | 83 ++++++++++++ .../glsl/math-functions.main.Fragment.glsl | 6 +- naga/tests/out/hlsl/math-functions.hlsl | 6 +- naga/tests/out/msl/math-functions.msl | 6 +- naga/tests/out/spv/math-functions.spvasm | 121 +++++++++--------- naga/tests/out/wgsl/math-functions.wgsl | 6 +- 6 files changed, 148 insertions(+), 80 deletions(-) diff --git a/naga/src/proc/constant_evaluator.rs b/naga/src/proc/constant_evaluator.rs index 344e2ddba2..deaa9c93c7 100644 --- a/naga/src/proc/constant_evaluator.rs +++ b/naga/src/proc/constant_evaluator.rs @@ -1233,6 +1233,9 @@ impl<'a> ConstantEvaluator<'a> { crate::MathFunction::ReverseBits => { component_wise_concrete_int!(self, span, [arg], |e| { Ok([e.reverse_bits()]) }) } + crate::MathFunction::FirstTrailingBit => { + component_wise_concrete_int(self, span, [arg], |ci| Ok(first_trailing_bit(ci))) + } crate::MathFunction::FirstLeadingBit => { component_wise_concrete_int(self, span, [arg], |ci| Ok(first_leading_bit(ci))) } @@ -2101,6 +2104,86 @@ impl<'a> ConstantEvaluator<'a> { } } +fn first_trailing_bit(concrete_int: ConcreteInt<1>) -> ConcreteInt<1> { + // NOTE: Bit indices for this built-in start at 0 at the "right" (or LSB). For example, a value + // of 1 means the least significant bit is set. Therefore, an input of `0x[80 00…]` would + // return a right-to-left bit index of 0. + let trailing_zeros_to_bit_idx = |e: u32| -> u32 { + match e { + idx @ 0..=31 => idx, + 32 => u32::MAX, + _ => unreachable!(), + } + }; + match concrete_int { + ConcreteInt::U32([e]) => ConcreteInt::U32([trailing_zeros_to_bit_idx(e.trailing_zeros())]), + ConcreteInt::I32([e]) => { + ConcreteInt::I32([trailing_zeros_to_bit_idx(e.trailing_zeros()) as i32]) + } + } +} + +#[test] +fn first_trailing_bit_smoke() { + assert_eq!( + first_trailing_bit(ConcreteInt::I32([0])), + ConcreteInt::I32([-1]) + ); + assert_eq!( + first_trailing_bit(ConcreteInt::I32([1])), + ConcreteInt::I32([0]) + ); + assert_eq!( + first_trailing_bit(ConcreteInt::I32([2])), + ConcreteInt::I32([1]) + ); + assert_eq!( + first_trailing_bit(ConcreteInt::I32([-1])), + ConcreteInt::I32([0]), + ); + assert_eq!( + first_trailing_bit(ConcreteInt::I32([i32::MIN])), + ConcreteInt::I32([31]), + ); + assert_eq!( + first_trailing_bit(ConcreteInt::I32([i32::MAX])), + ConcreteInt::I32([0]), + ); + for idx in 0..32 { + assert_eq!( + first_trailing_bit(ConcreteInt::I32([1 << idx])), + ConcreteInt::I32([idx]) + ) + } + + assert_eq!( + first_trailing_bit(ConcreteInt::U32([0])), + ConcreteInt::U32([u32::MAX]) + ); + assert_eq!( + first_trailing_bit(ConcreteInt::U32([1])), + ConcreteInt::U32([0]) + ); + assert_eq!( + first_trailing_bit(ConcreteInt::U32([2])), + ConcreteInt::U32([1]) + ); + assert_eq!( + first_trailing_bit(ConcreteInt::U32([1 << 31])), + ConcreteInt::U32([31]), + ); + assert_eq!( + first_trailing_bit(ConcreteInt::U32([u32::MAX])), + ConcreteInt::U32([0]), + ); + for idx in 0..32 { + assert_eq!( + first_trailing_bit(ConcreteInt::U32([1 << idx])), + ConcreteInt::U32([idx]) + ) + } +} + fn first_leading_bit(concrete_int: ConcreteInt<1>) -> ConcreteInt<1> { // NOTE: Bit indices for this built-in start at 0 at the "right" (or LSB). For example, 1 means // the least significant bit is set. Therefore, an input of 1 would return a right-to-left bit diff --git a/naga/tests/out/glsl/math-functions.main.Fragment.glsl b/naga/tests/out/glsl/math-functions.main.Fragment.glsl index c10dddff0f..4ab85269e1 100644 --- a/naga/tests/out/glsl/math-functions.main.Fragment.glsl +++ b/naga/tests/out/glsl/math-functions.main.Fragment.glsl @@ -67,10 +67,8 @@ void main() { int const_dot = ( + ivec2(0).x * ivec2(0).x + ivec2(0).y * ivec2(0).y); ivec2 flb_b = ivec2(-1, -1); uvec2 flb_c = uvec2(0u, 0u); - int ftb_a = findLSB(-1); - uint ftb_b = uint(findLSB(1u)); - ivec2 ftb_c = findLSB(ivec2(-1)); - uvec2 ftb_d = uvec2(findLSB(uvec2(1u))); + ivec2 ftb_c = ivec2(0, 0); + uvec2 ftb_d = uvec2(0u, 0u); uvec2 ctz_e = uvec2(32u, 32u); ivec2 ctz_f = ivec2(32, 32); uvec2 ctz_g = uvec2(0u, 0u); diff --git a/naga/tests/out/hlsl/math-functions.hlsl b/naga/tests/out/hlsl/math-functions.hlsl index 14d1e9e188..a02b2b1280 100644 --- a/naga/tests/out/hlsl/math-functions.hlsl +++ b/naga/tests/out/hlsl/math-functions.hlsl @@ -81,10 +81,8 @@ void main() int const_dot = dot(ZeroValueint2(), ZeroValueint2()); int2 flb_b = int2(-1, -1); uint2 flb_c = uint2(0u, 0u); - int ftb_a = asint(firstbitlow(-1)); - uint ftb_b = firstbitlow(1u); - int2 ftb_c = asint(firstbitlow((-1).xx)); - uint2 ftb_d = firstbitlow((1u).xx); + int2 ftb_c = int2(0, 0); + uint2 ftb_d = uint2(0u, 0u); uint2 ctz_e = uint2(32u, 32u); int2 ctz_f = int2(32, 32); uint2 ctz_g = uint2(0u, 0u); diff --git a/naga/tests/out/msl/math-functions.msl b/naga/tests/out/msl/math-functions.msl index 271472978a..559002c39b 100644 --- a/naga/tests/out/msl/math-functions.msl +++ b/naga/tests/out/msl/math-functions.msl @@ -69,10 +69,8 @@ fragment void main_( int const_dot = ( + metal::int2 {}.x * metal::int2 {}.x + metal::int2 {}.y * metal::int2 {}.y); metal::int2 flb_b = metal::int2(-1, -1); metal::uint2 flb_c = metal::uint2(0u, 0u); - int ftb_a = (((metal::ctz(-1) + 1) % 33) - 1); - uint ftb_b = (((metal::ctz(1u) + 1) % 33) - 1); - metal::int2 ftb_c = (((metal::ctz(metal::int2(-1)) + 1) % 33) - 1); - metal::uint2 ftb_d = (((metal::ctz(metal::uint2(1u)) + 1) % 33) - 1); + metal::int2 ftb_c = metal::int2(0, 0); + metal::uint2 ftb_d = metal::uint2(0u, 0u); metal::uint2 ctz_e = metal::uint2(32u, 32u); metal::int2 ctz_f = metal::int2(32, 32); metal::uint2 ctz_g = metal::uint2(0u, 0u); diff --git a/naga/tests/out/spv/math-functions.spvasm b/naga/tests/out/spv/math-functions.spvasm index 2207934cc9..366857f91f 100644 --- a/naga/tests/out/spv/math-functions.spvasm +++ b/naga/tests/out/spv/math-functions.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 94 +; Bound: 87 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 @@ -44,71 +44,64 @@ OpMemberDecorate %15 1 Offset 16 %28 = OpConstantComposite %7 %22 %22 %29 = OpConstant %9 0 %30 = OpConstantComposite %8 %29 %29 -%31 = OpConstant %9 1 -%32 = OpConstantComposite %7 %22 %22 -%33 = OpConstantComposite %8 %31 %31 -%34 = OpConstant %9 32 -%35 = OpConstant %6 32 -%36 = OpConstant %6 0 -%37 = OpConstantComposite %8 %34 %34 -%38 = OpConstantComposite %7 %35 %35 -%39 = OpConstantComposite %7 %36 %36 -%40 = OpConstant %9 31 -%41 = OpConstantComposite %8 %40 %40 -%42 = OpConstant %6 2 -%43 = OpConstant %4 2.0 -%44 = OpConstantComposite %10 %19 %43 -%45 = OpConstant %6 3 -%46 = OpConstant %6 4 -%47 = OpConstantComposite %7 %45 %46 -%48 = OpConstant %4 1.5 -%49 = OpConstantComposite %10 %48 %48 -%50 = OpConstantComposite %3 %48 %48 %48 %48 -%57 = OpConstantComposite %3 %19 %19 %19 %19 -%60 = OpConstantNull %6 +%31 = OpConstant %6 0 +%32 = OpConstantComposite %7 %31 %31 +%33 = OpConstant %9 32 +%34 = OpConstant %6 32 +%35 = OpConstantComposite %8 %33 %33 +%36 = OpConstantComposite %7 %34 %34 +%37 = OpConstant %9 31 +%38 = OpConstantComposite %8 %37 %37 +%39 = OpConstant %6 2 +%40 = OpConstant %4 2.0 +%41 = OpConstantComposite %10 %19 %40 +%42 = OpConstant %6 3 +%43 = OpConstant %6 4 +%44 = OpConstantComposite %7 %42 %43 +%45 = OpConstant %4 1.5 +%46 = OpConstantComposite %10 %45 %45 +%47 = OpConstantComposite %3 %45 %45 %45 %45 +%54 = OpConstantComposite %3 %19 %19 %19 %19 +%57 = OpConstantNull %6 %17 = OpFunction %2 None %18 %16 = OpLabel -OpBranch %51 -%51 = OpLabel -%52 = OpExtInst %4 %1 Degrees %19 -%53 = OpExtInst %4 %1 Radians %19 -%54 = OpExtInst %3 %1 Degrees %21 -%55 = OpExtInst %3 %1 Radians %21 -%56 = OpExtInst %3 %1 FClamp %21 %21 %57 -%58 = OpExtInst %3 %1 Refract %21 %21 %19 -%61 = OpCompositeExtract %6 %26 0 -%62 = OpCompositeExtract %6 %26 0 -%63 = OpIMul %6 %61 %62 -%64 = OpIAdd %6 %60 %63 -%65 = OpCompositeExtract %6 %26 1 -%66 = OpCompositeExtract %6 %26 1 -%67 = OpIMul %6 %65 %66 -%59 = OpIAdd %6 %64 %67 -%68 = OpExtInst %6 %1 FindILsb %22 -%69 = OpExtInst %9 %1 FindILsb %31 -%70 = OpExtInst %7 %1 FindILsb %32 -%71 = OpExtInst %8 %1 FindILsb %33 -%72 = OpExtInst %4 %1 Ldexp %19 %42 -%73 = OpExtInst %10 %1 Ldexp %44 %47 -%74 = OpExtInst %11 %1 ModfStruct %48 -%75 = OpExtInst %11 %1 ModfStruct %48 -%76 = OpCompositeExtract %4 %75 0 -%77 = OpExtInst %11 %1 ModfStruct %48 +OpBranch %48 +%48 = OpLabel +%49 = OpExtInst %4 %1 Degrees %19 +%50 = OpExtInst %4 %1 Radians %19 +%51 = OpExtInst %3 %1 Degrees %21 +%52 = OpExtInst %3 %1 Radians %21 +%53 = OpExtInst %3 %1 FClamp %21 %21 %54 +%55 = OpExtInst %3 %1 Refract %21 %21 %19 +%58 = OpCompositeExtract %6 %26 0 +%59 = OpCompositeExtract %6 %26 0 +%60 = OpIMul %6 %58 %59 +%61 = OpIAdd %6 %57 %60 +%62 = OpCompositeExtract %6 %26 1 +%63 = OpCompositeExtract %6 %26 1 +%64 = OpIMul %6 %62 %63 +%56 = OpIAdd %6 %61 %64 +%65 = OpExtInst %4 %1 Ldexp %19 %39 +%66 = OpExtInst %10 %1 Ldexp %41 %44 +%67 = OpExtInst %11 %1 ModfStruct %45 +%68 = OpExtInst %11 %1 ModfStruct %45 +%69 = OpCompositeExtract %4 %68 0 +%70 = OpExtInst %11 %1 ModfStruct %45 +%71 = OpCompositeExtract %4 %70 1 +%72 = OpExtInst %12 %1 ModfStruct %46 +%73 = OpExtInst %13 %1 ModfStruct %47 +%74 = OpCompositeExtract %3 %73 1 +%75 = OpCompositeExtract %4 %74 0 +%76 = OpExtInst %12 %1 ModfStruct %46 +%77 = OpCompositeExtract %10 %76 0 %78 = OpCompositeExtract %4 %77 1 -%79 = OpExtInst %12 %1 ModfStruct %49 -%80 = OpExtInst %13 %1 ModfStruct %50 -%81 = OpCompositeExtract %3 %80 1 -%82 = OpCompositeExtract %4 %81 0 -%83 = OpExtInst %12 %1 ModfStruct %49 -%84 = OpCompositeExtract %10 %83 0 -%85 = OpCompositeExtract %4 %84 1 -%86 = OpExtInst %14 %1 FrexpStruct %48 -%87 = OpExtInst %14 %1 FrexpStruct %48 -%88 = OpCompositeExtract %4 %87 0 -%89 = OpExtInst %14 %1 FrexpStruct %48 -%90 = OpCompositeExtract %6 %89 1 -%91 = OpExtInst %15 %1 FrexpStruct %50 -%92 = OpCompositeExtract %5 %91 1 -%93 = OpCompositeExtract %6 %92 0 +%79 = OpExtInst %14 %1 FrexpStruct %45 +%80 = OpExtInst %14 %1 FrexpStruct %45 +%81 = OpCompositeExtract %4 %80 0 +%82 = OpExtInst %14 %1 FrexpStruct %45 +%83 = OpCompositeExtract %6 %82 1 +%84 = OpExtInst %15 %1 FrexpStruct %47 +%85 = OpCompositeExtract %5 %84 1 +%86 = OpCompositeExtract %6 %85 0 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/math-functions.wgsl b/naga/tests/out/wgsl/math-functions.wgsl index c97ce6a4a2..2271bb9cb0 100644 --- a/naga/tests/out/wgsl/math-functions.wgsl +++ b/naga/tests/out/wgsl/math-functions.wgsl @@ -12,10 +12,8 @@ fn main() { let const_dot = dot(vec2(), vec2()); let flb_b = vec2(-1i, -1i); let flb_c = vec2(0u, 0u); - let ftb_a = firstTrailingBit(-1i); - let ftb_b = firstTrailingBit(1u); - let ftb_c = firstTrailingBit(vec2(-1i)); - let ftb_d = firstTrailingBit(vec2(1u)); + let ftb_c = vec2(0i, 0i); + let ftb_d = vec2(0u, 0u); let ctz_e = vec2(32u, 32u); let ctz_f = vec2(32i, 32i); let ctz_g = vec2(0u, 0u); From a1e6c860397ec2680a6530e9d1e6a9bda480e50a Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Tue, 23 Jul 2024 12:45:48 -0400 Subject: [PATCH 7/7] docs(CHANGELOG): add entry for const. eval of `first{Leading,Trailing}Bit` --- CHANGELOG.md | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index c9eccafcda..9567c74f0f 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -39,6 +39,12 @@ Bottom level categories: ## Unreleased +### New Features + +#### Naga + +* Support constant evaluation for `firstLeadingBit` and `firstTrailingBit` numeric built-ins in WGSL. Front-ends that translate to these built-ins also benefit from constant evaluation. By @ErichDonGubler in [#5101](https://github.com/gfx-rs/wgpu/pull/5101). + ### Bug Fixes #### General