Skip to content

Commit a7a2f16

Browse files
ErichDonGublerErichDonGubler
committed
feat(const_eval): impl. firstLeadingBit 
1 parent 33f6789 commit a7a2f16

File tree

6 files changed

+97
-87
lines changed

6 files changed

+97
-87
lines changed

‎naga/src/proc/constant_evaluator.rs

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1228,6 +1228,28 @@ impl<'a> ConstantEvaluator<'a> {
12281228
crate::MathFunction::ReverseBits => {
12291229
component_wise_concrete_int!(self, span, [arg], |e| { Ok([e.reverse_bits()]) })
12301230
}
1231+
crate::MathFunction::FirstLeadingBit => {
1232+
component_wise_concrete_int(self, span, [arg], |concrete_int| match concrete_int {
1233+
ConcreteInt::I32([e]) => Ok(ConcreteInt::I32([{
1234+
if matches!(e, 0 | -1) {
1235+
-1
1236+
} else if e.is_positive() {
1237+
32 /* bits */ - i32::try_from(e.leading_zeros())
1238+
.expect("bit count overflowed 32 bits, somehow!?")
1239+
} else {
1240+
32 /* bits */ - i32::try_from(e.leading_ones())
1241+
.expect("bit count overflowed 32 bits, somehow!?")
1242+
}
1243+
}])),
1244+
ConcreteInt::U32([e]) => Ok(ConcreteInt::U32([{
1245+
if e == 0 {
1246+
u32::MAX
1247+
} else {
1248+
e.leading_zeros()
1249+
}
1250+
}])),
1251+
})
1252+
}
12311253

12321254
fun => Err(ConstantEvaluatorError::NotImplemented(format!(
12331255
"{fun:?} built-in function"

‎naga/tests/out/glsl/math-functions.main.Fragment.glsl

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -65,10 +65,8 @@ void main() {
6565
ivec4 sign_b = ivec4(-1, -1, -1, -1);
6666
vec4 sign_d = vec4(-1.0, -1.0, -1.0, -1.0);
6767
int const_dot = ( + ivec2(0).x * ivec2(0).x + ivec2(0).y * ivec2(0).y);
68-
uint first_leading_bit_abs = uint(findMSB(0u));
69-
int flb_a = findMSB(-1);
70-
ivec2 flb_b = findMSB(ivec2(-1));
71-
uvec2 flb_c = uvec2(findMSB(uvec2(1u)));
68+
ivec2 flb_b = ivec2(-1, -1);
69+
uvec2 flb_c = uvec2(31u, 31u);
7270
int ftb_a = findLSB(-1);
7371
uint ftb_b = uint(findLSB(1u));
7472
ivec2 ftb_c = findLSB(ivec2(-1));

‎naga/tests/out/hlsl/math-functions.hlsl

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -75,10 +75,8 @@ void main()
7575
int4 sign_b = int4(-1, -1, -1, -1);
7676
float4 sign_d = float4(-1.0, -1.0, -1.0, -1.0);
7777
int const_dot = dot((int2)0, (int2)0);
78-
uint first_leading_bit_abs = firstbithigh(0u);
79-
int flb_a = asint(firstbithigh(-1));
80-
int2 flb_b = asint(firstbithigh((-1).xx));
81-
uint2 flb_c = firstbithigh((1u).xx);
78+
int2 flb_b = int2(-1, -1);
79+
uint2 flb_c = uint2(31u, 31u);
8280
int ftb_a = asint(firstbitlow(-1));
8381
uint ftb_b = firstbitlow(1u);
8482
int2 ftb_c = asint(firstbitlow((-1).xx));

‎naga/tests/out/msl/math-functions.msl

Lines changed: 2 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -67,12 +67,8 @@ fragment void main_(
6767
metal::int4 sign_b = metal::int4(-1, -1, -1, -1);
6868
metal::float4 sign_d = metal::float4(-1.0, -1.0, -1.0, -1.0);
6969
int const_dot = ( + metal::int2 {}.x * metal::int2 {}.x + metal::int2 {}.y * metal::int2 {}.y);
70-
uint first_leading_bit_abs = metal::select(31 - metal::clz(0u), uint(-1), 0u == 0 || 0u == -1);
71-
int flb_a = metal::select(31 - metal::clz(metal::select(-1, ~-1, -1 < 0)), int(-1), -1 == 0 || -1 == -1);
72-
metal::int2 _e29 = metal::int2(-1);
73-
metal::int2 flb_b = metal::select(31 - metal::clz(metal::select(_e29, ~_e29, _e29 < 0)), int2(-1), _e29 == 0 || _e29 == -1);
74-
metal::uint2 _e32 = metal::uint2(1u);
75-
metal::uint2 flb_c = metal::select(31 - metal::clz(_e32), uint2(-1), _e32 == 0 || _e32 == -1);
70+
metal::int2 flb_b = metal::int2(-1, -1);
71+
metal::uint2 flb_c = metal::uint2(31u, 31u);
7672
int ftb_a = (((metal::ctz(-1) + 1) % 33) - 1);
7773
uint ftb_b = (((metal::ctz(1u) + 1) % 33) - 1);
7874
metal::int2 ftb_c = (((metal::ctz(metal::int2(-1)) + 1) % 33) - 1);

‎naga/tests/out/spv/math-functions.spvasm

Lines changed: 67 additions & 69 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
; SPIR-V
22
; Version: 1.1
33
; Generator: rspirv
4-
; Bound: 96
4+
; Bound: 94
55
OpCapability Shader
66
%1 = OpExtInstImport "GLSL.std.450"
77
OpMemoryModel Logical GLSL450
@@ -40,77 +40,75 @@ OpMemberDecorate %15 1 Offset 16
4040
%24 = OpConstant %4 -1.0
4141
%25 = OpConstantComposite %3 %24 %24 %24 %24
4242
%26 = OpConstantNull %7
43-
%27 = OpConstant %9 0
43+
%27 = OpConstant %9 4294967295
4444
%28 = OpConstantComposite %7 %22 %22
45-
%29 = OpConstant %9 1
45+
%29 = OpConstant %9 31
4646
%30 = OpConstantComposite %8 %29 %29
47-
%31 = OpConstant %9 32
48-
%32 = OpConstant %6 32
49-
%33 = OpConstant %6 0
50-
%34 = OpConstantComposite %8 %31 %31
51-
%35 = OpConstantComposite %7 %32 %32
52-
%36 = OpConstantComposite %8 %27 %27
53-
%37 = OpConstantComposite %7 %33 %33
54-
%38 = OpConstant %9 31
55-
%39 = OpConstantComposite %8 %38 %38
56-
%40 = OpConstant %6 2
57-
%41 = OpConstant %4 2.0
58-
%42 = OpConstantComposite %10 %19 %41
59-
%43 = OpConstant %6 3
60-
%44 = OpConstant %6 4
61-
%45 = OpConstantComposite %7 %43 %44
62-
%46 = OpConstant %4 1.5
63-
%47 = OpConstantComposite %10 %46 %46
64-
%48 = OpConstantComposite %3 %46 %46 %46 %46
65-
%55 = OpConstantComposite %3 %19 %19 %19 %19
66-
%58 = OpConstantNull %6
47+
%31 = OpConstant %9 1
48+
%32 = OpConstantComposite %7 %22 %22
49+
%33 = OpConstantComposite %8 %31 %31
50+
%34 = OpConstant %9 32
51+
%35 = OpConstant %6 32
52+
%36 = OpConstant %9 0
53+
%37 = OpConstant %6 0
54+
%38 = OpConstantComposite %8 %34 %34
55+
%39 = OpConstantComposite %7 %35 %35
56+
%40 = OpConstantComposite %8 %36 %36
57+
%41 = OpConstantComposite %7 %37 %37
58+
%42 = OpConstant %6 2
59+
%43 = OpConstant %4 2.0
60+
%44 = OpConstantComposite %10 %19 %43
61+
%45 = OpConstant %6 3
62+
%46 = OpConstant %6 4
63+
%47 = OpConstantComposite %7 %45 %46
64+
%48 = OpConstant %4 1.5
65+
%49 = OpConstantComposite %10 %48 %48
66+
%50 = OpConstantComposite %3 %48 %48 %48 %48
67+
%57 = OpConstantComposite %3 %19 %19 %19 %19
68+
%60 = OpConstantNull %6
6769
%17 = OpFunction %2 None %18
6870
%16 = OpLabel
69-
OpBranch %49
70-
%49 = OpLabel
71-
%50 = OpExtInst %4 %1 Degrees %19
72-
%51 = OpExtInst %4 %1 Radians %19
73-
%52 = OpExtInst %3 %1 Degrees %21
74-
%53 = OpExtInst %3 %1 Radians %21
75-
%54 = OpExtInst %3 %1 FClamp %21 %21 %55
76-
%56 = OpExtInst %3 %1 Refract %21 %21 %19
77-
%59 = OpCompositeExtract %6 %26 0
78-
%60 = OpCompositeExtract %6 %26 0
79-
%61 = OpIMul %6 %59 %60
80-
%62 = OpIAdd %6 %58 %61
81-
%63 = OpCompositeExtract %6 %26 1
82-
%64 = OpCompositeExtract %6 %26 1
83-
%65 = OpIMul %6 %63 %64
84-
%57 = OpIAdd %6 %62 %65
85-
%66 = OpExtInst %9 %1 FindUMsb %27
86-
%67 = OpExtInst %6 %1 FindSMsb %22
87-
%68 = OpExtInst %7 %1 FindSMsb %28
88-
%69 = OpExtInst %8 %1 FindUMsb %30
89-
%70 = OpExtInst %6 %1 FindILsb %22
90-
%71 = OpExtInst %9 %1 FindILsb %29
91-
%72 = OpExtInst %7 %1 FindILsb %28
92-
%73 = OpExtInst %8 %1 FindILsb %30
93-
%74 = OpExtInst %4 %1 Ldexp %19 %40
94-
%75 = OpExtInst %10 %1 Ldexp %42 %45
95-
%76 = OpExtInst %11 %1 ModfStruct %46
96-
%77 = OpExtInst %11 %1 ModfStruct %46
97-
%78 = OpCompositeExtract %4 %77 0
98-
%79 = OpExtInst %11 %1 ModfStruct %46
99-
%80 = OpCompositeExtract %4 %79 1
100-
%81 = OpExtInst %12 %1 ModfStruct %47
101-
%82 = OpExtInst %13 %1 ModfStruct %48
102-
%83 = OpCompositeExtract %3 %82 1
103-
%84 = OpCompositeExtract %4 %83 0
104-
%85 = OpExtInst %12 %1 ModfStruct %47
105-
%86 = OpCompositeExtract %10 %85 0
106-
%87 = OpCompositeExtract %4 %86 1
107-
%88 = OpExtInst %14 %1 FrexpStruct %46
108-
%89 = OpExtInst %14 %1 FrexpStruct %46
109-
%90 = OpCompositeExtract %4 %89 0
110-
%91 = OpExtInst %14 %1 FrexpStruct %46
111-
%92 = OpCompositeExtract %6 %91 1
112-
%93 = OpExtInst %15 %1 FrexpStruct %48
113-
%94 = OpCompositeExtract %5 %93 1
114-
%95 = OpCompositeExtract %6 %94 0
71+
OpBranch %51
72+
%51 = OpLabel
73+
%52 = OpExtInst %4 %1 Degrees %19
74+
%53 = OpExtInst %4 %1 Radians %19
75+
%54 = OpExtInst %3 %1 Degrees %21
76+
%55 = OpExtInst %3 %1 Radians %21
77+
%56 = OpExtInst %3 %1 FClamp %21 %21 %57
78+
%58 = OpExtInst %3 %1 Refract %21 %21 %19
79+
%61 = OpCompositeExtract %6 %26 0
80+
%62 = OpCompositeExtract %6 %26 0
81+
%63 = OpIMul %6 %61 %62
82+
%64 = OpIAdd %6 %60 %63
83+
%65 = OpCompositeExtract %6 %26 1
84+
%66 = OpCompositeExtract %6 %26 1
85+
%67 = OpIMul %6 %65 %66
86+
%59 = OpIAdd %6 %64 %67
87+
%68 = OpExtInst %6 %1 FindILsb %22
88+
%69 = OpExtInst %9 %1 FindILsb %31
89+
%70 = OpExtInst %7 %1 FindILsb %32
90+
%71 = OpExtInst %8 %1 FindILsb %33
91+
%72 = OpExtInst %4 %1 Ldexp %19 %42
92+
%73 = OpExtInst %10 %1 Ldexp %44 %47
93+
%74 = OpExtInst %11 %1 ModfStruct %48
94+
%75 = OpExtInst %11 %1 ModfStruct %48
95+
%76 = OpCompositeExtract %4 %75 0
96+
%77 = OpExtInst %11 %1 ModfStruct %48
97+
%78 = OpCompositeExtract %4 %77 1
98+
%79 = OpExtInst %12 %1 ModfStruct %49
99+
%80 = OpExtInst %13 %1 ModfStruct %50
100+
%81 = OpCompositeExtract %3 %80 1
101+
%82 = OpCompositeExtract %4 %81 0
102+
%83 = OpExtInst %12 %1 ModfStruct %49
103+
%84 = OpCompositeExtract %10 %83 0
104+
%85 = OpCompositeExtract %4 %84 1
105+
%86 = OpExtInst %14 %1 FrexpStruct %48
106+
%87 = OpExtInst %14 %1 FrexpStruct %48
107+
%88 = OpCompositeExtract %4 %87 0
108+
%89 = OpExtInst %14 %1 FrexpStruct %48
109+
%90 = OpCompositeExtract %6 %89 1
110+
%91 = OpExtInst %15 %1 FrexpStruct %50
111+
%92 = OpCompositeExtract %5 %91 1
112+
%93 = OpCompositeExtract %6 %92 0
115113
OpReturn
116114
OpFunctionEnd

‎naga/tests/out/wgsl/math-functions.wgsl

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -10,10 +10,8 @@ fn main() {
1010
let sign_b = vec4<i32>(-1i, -1i, -1i, -1i);
1111
let sign_d = vec4<f32>(-1f, -1f, -1f, -1f);
1212
let const_dot = dot(vec2<i32>(), vec2<i32>());
13-
let first_leading_bit_abs = firstLeadingBit(0u);
14-
let flb_a = firstLeadingBit(-1i);
15-
let flb_b = firstLeadingBit(vec2(-1i));
16-
let flb_c = firstLeadingBit(vec2(1u));
13+
let flb_b = vec2<i32>(-1i, -1i);
14+
let flb_c = vec2<u32>(31u, 31u);
1715
let ftb_a = firstTrailingBit(-1i);
1816
let ftb_b = firstTrailingBit(1u);
1917
let ftb_c = firstTrailingBit(vec2(-1i));

0 commit comments

Comments
 (0)