nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-10-30 14:01:39 +00:00
Code Issues Packages Projects Releases Wiki Activity

feat(const_eval): impl. firstTrailingBit

Browse Source
This commit is contained in:
Erich Gubler 2024-01-26 13:06:20 -05:00
parent 2f7c87f7af
commit a220fcfc57
6 changed files with 148 additions and 80 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

83
naga/src/proc/constant_evaluator.rs
View File

@ -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

6
naga/tests/out/glsl/math-functions.main.Fragment.glsl
View File

@ -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);

6
naga/tests/out/hlsl/math-functions.hlsl
View File

@ -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);

6
naga/tests/out/msl/math-functions.msl
View File

@ -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);

121
naga/tests/out/spv/math-functions.spvasm
View File

@ -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

6
naga/tests/out/wgsl/math-functions.wgsl
View File

@ -12,10 +12,8 @@ fn main() {
let const_dot = dot(vec2<i32>(), vec2<i32>());
let flb_b = vec2<i32>(-1i, -1i);
let flb_c = vec2<u32>(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<i32>(0i, 0i);
let ftb_d = vec2<u32>(0u, 0u);
let ctz_e = vec2<u32>(32u, 32u);
let ctz_f = vec2<i32>(32i, 32i);
let ctz_g = vec2<u32>(0u, 0u);