Skip to content

Commit 2b36f65

Browse files
ErichDonGublerErichDonGubler
committed
feat(const_eval): impl. firstTrailingBit 
1 parent 09c580b commit 2b36f65

File tree

6 files changed

+80
-79
lines changed

6 files changed

+80
-79
lines changed

‎naga/src/proc/constant_evaluator.rs

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1039,6 +1039,20 @@ impl<'a> ConstantEvaluator<'a> {
10391039
crate::MathFunction::ReverseBits => {
10401040
component_wise_concrete_int!(self, span, [arg], |e| { Ok([e.reverse_bits()]) })
10411041
}
1042+
crate::MathFunction::FirstTrailingBit => {
1043+
component_wise_concrete_int(self, span, [arg], |concrete_int| match concrete_int {
1044+
ConcreteInt::U32([e]) => Ok(ConcreteInt::U32([{
1045+
(32 /* bits */ - e.trailing_zeros())
1046+
.checked_sub(1)
1047+
.unwrap_or(u32::MAX)
1048+
}])),
1049+
ConcreteInt::I32([e]) => Ok(ConcreteInt::I32([{
1050+
(32 /* bits */ - i32::try_from(e.trailing_zeros())
1051+
.expect("bit count overflowed 32 bits, somehow!?"))
1052+
- 1
1053+
}])),
1054+
})
1055+
}
10421056
crate::MathFunction::FirstLeadingBit => {
10431057
component_wise_concrete_int(self, span, [arg], |concrete_int| match concrete_int {
10441058
ConcreteInt::I32([e]) => Ok(ConcreteInt::I32([{

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

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -67,10 +67,8 @@ void main() {
6767
int const_dot = ( + ivec2(0).x * ivec2(0).x + ivec2(0).y * ivec2(0).y);
6868
ivec2 flb_b = ivec2(-1, -1);
6969
uvec2 flb_c = uvec2(31u, 31u);
70-
int ftb_a = findLSB(-1);
71-
uint ftb_b = uint(findLSB(1u));
72-
ivec2 ftb_c = findLSB(ivec2(-1));
73-
uvec2 ftb_d = uvec2(findLSB(uvec2(1u)));
70+
ivec2 ftb_c = ivec2(31, 31);
71+
uvec2 ftb_d = uvec2(31u, 31u);
7472
uvec2 ctz_e = uvec2(32u, 32u);
7573
ivec2 ctz_f = ivec2(32, 32);
7674
uvec2 ctz_g = uvec2(0u, 0u);

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

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -77,10 +77,8 @@ void main()
7777
int const_dot = dot((int2)0, (int2)0);
7878
int2 flb_b = int2(-1, -1);
7979
uint2 flb_c = uint2(31u, 31u);
80-
int ftb_a = asint(firstbitlow(-1));
81-
uint ftb_b = firstbitlow(1u);
82-
int2 ftb_c = asint(firstbitlow((-1).xx));
83-
uint2 ftb_d = firstbitlow((1u).xx);
80+
int2 ftb_c = int2(31, 31);
81+
uint2 ftb_d = uint2(31u, 31u);
8482
uint2 ctz_e = uint2(32u, 32u);
8583
int2 ctz_f = int2(32, 32);
8684
uint2 ctz_g = uint2(0u, 0u);

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

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -69,10 +69,8 @@ fragment void main_(
6969
int const_dot = ( + metal::int2 {}.x * metal::int2 {}.x + metal::int2 {}.y * metal::int2 {}.y);
7070
metal::int2 flb_b = metal::int2(-1, -1);
7171
metal::uint2 flb_c = metal::uint2(31u, 31u);
72-
int ftb_a = (((metal::ctz(-1) + 1) % 33) - 1);
73-
uint ftb_b = (((metal::ctz(1u) + 1) % 33) - 1);
74-
metal::int2 ftb_c = (((metal::ctz(metal::int2(-1)) + 1) % 33) - 1);
75-
metal::uint2 ftb_d = (((metal::ctz(metal::uint2(1u)) + 1) % 33) - 1);
72+
metal::int2 ftb_c = metal::int2(31, 31);
73+
metal::uint2 ftb_d = metal::uint2(31u, 31u);
7674
metal::uint2 ctz_e = metal::uint2(32u, 32u);
7775
metal::int2 ctz_f = metal::int2(32, 32);
7876
metal::uint2 ctz_g = metal::uint2(0u, 0u);

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

Lines changed: 58 additions & 63 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: 94
4+
; Bound: 89
55
OpCapability Shader
66
%1 = OpExtInstImport "GLSL.std.450"
77
OpMemoryModel Logical GLSL450
@@ -44,71 +44,66 @@ OpMemberDecorate %15 1 Offset 16
4444
%28 = OpConstantComposite %7 %22 %22
4545
%29 = OpConstant %9 31
4646
%30 = OpConstantComposite %8 %29 %29
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
47+
%31 = OpConstant %6 31
48+
%32 = OpConstantComposite %7 %31 %31
49+
%33 = OpConstant %9 32
50+
%34 = OpConstant %6 32
51+
%35 = OpConstant %9 0
52+
%36 = OpConstant %6 0
53+
%37 = OpConstantComposite %8 %33 %33
54+
%38 = OpConstantComposite %7 %34 %34
55+
%39 = OpConstantComposite %8 %35 %35
56+
%40 = OpConstantComposite %7 %36 %36
57+
%41 = OpConstant %6 2
58+
%42 = OpConstant %4 2.0
59+
%43 = OpConstantComposite %10 %19 %42
60+
%44 = OpConstant %6 3
61+
%45 = OpConstant %6 4
62+
%46 = OpConstantComposite %7 %44 %45
63+
%47 = OpConstant %4 1.5
64+
%48 = OpConstantComposite %10 %47 %47
65+
%49 = OpConstantComposite %3 %47 %47 %47 %47
66+
%56 = OpConstantComposite %3 %19 %19 %19 %19
67+
%59 = OpConstantNull %6
6968
%17 = OpFunction %2 None %18
7069
%16 = OpLabel
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
70+
OpBranch %50
71+
%50 = OpLabel
72+
%51 = OpExtInst %4 %1 Degrees %19
73+
%52 = OpExtInst %4 %1 Radians %19
74+
%53 = OpExtInst %3 %1 Degrees %21
75+
%54 = OpExtInst %3 %1 Radians %21
76+
%55 = OpExtInst %3 %1 FClamp %21 %21 %56
77+
%57 = OpExtInst %3 %1 Refract %21 %21 %19
78+
%60 = OpCompositeExtract %6 %26 0
7979
%61 = OpCompositeExtract %6 %26 0
80-
%62 = OpCompositeExtract %6 %26 0
81-
%63 = OpIMul %6 %61 %62
82-
%64 = OpIAdd %6 %60 %63
80+
%62 = OpIMul %6 %60 %61
81+
%63 = OpIAdd %6 %59 %62
82+
%64 = OpCompositeExtract %6 %26 1
8383
%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
84+
%66 = OpIMul %6 %64 %65
85+
%58 = OpIAdd %6 %63 %66
86+
%67 = OpExtInst %4 %1 Ldexp %19 %41
87+
%68 = OpExtInst %10 %1 Ldexp %43 %46
88+
%69 = OpExtInst %11 %1 ModfStruct %47
89+
%70 = OpExtInst %11 %1 ModfStruct %47
90+
%71 = OpCompositeExtract %4 %70 0
91+
%72 = OpExtInst %11 %1 ModfStruct %47
92+
%73 = OpCompositeExtract %4 %72 1
93+
%74 = OpExtInst %12 %1 ModfStruct %48
94+
%75 = OpExtInst %13 %1 ModfStruct %49
95+
%76 = OpCompositeExtract %3 %75 1
96+
%77 = OpCompositeExtract %4 %76 0
97+
%78 = OpExtInst %12 %1 ModfStruct %48
98+
%79 = OpCompositeExtract %10 %78 0
99+
%80 = OpCompositeExtract %4 %79 1
100+
%81 = OpExtInst %14 %1 FrexpStruct %47
101+
%82 = OpExtInst %14 %1 FrexpStruct %47
102+
%83 = OpCompositeExtract %4 %82 0
103+
%84 = OpExtInst %14 %1 FrexpStruct %47
104+
%85 = OpCompositeExtract %6 %84 1
105+
%86 = OpExtInst %15 %1 FrexpStruct %49
106+
%87 = OpCompositeExtract %5 %86 1
107+
%88 = OpCompositeExtract %6 %87 0
113108
OpReturn
114109
OpFunctionEnd

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

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -12,10 +12,8 @@ fn main() {
1212
let const_dot = dot(vec2<i32>(), vec2<i32>());
1313
let flb_b = vec2<i32>(-1i, -1i);
1414
let flb_c = vec2<u32>(31u, 31u);
15-
let ftb_a = firstTrailingBit(-1i);
16-
let ftb_b = firstTrailingBit(1u);
17-
let ftb_c = firstTrailingBit(vec2(-1i));
18-
let ftb_d = firstTrailingBit(vec2(1u));
15+
let ftb_c = vec2<i32>(31i, 31i);
16+
let ftb_d = vec2<u32>(31u, 31u);
1917
let ctz_e = vec2<u32>(32u, 32u);
2018
let ctz_f = vec2<i32>(32i, 32i);
2119
let ctz_g = vec2<u32>(0u, 0u);

0 commit comments

Comments
 (0)