mirror of
https://github.com/gfx-rs/wgpu.git
synced 2024-11-25 16:24:24 +00:00
feat(const_eval): impl. countTrailingZeros
This commit is contained in:
Erich Gubler
Teodor Tanasoaia
parent
2ccc4f49ee
commit
66fd1872ab
@ -70,6 +70,7 @@ Bottom level categories:
|
||||
- `ceil`
|
||||
- `countLeadingZeros`
|
||||
- `countOneBits`
|
||||
- `countTrailingZeros`
|
||||
- `floor`
|
||||
- Eager release of GPU resources comes from device.trackers. By @bradwerth in [#5075](https://github.com/gfx-rs/wgpu/pull/5075)
|
||||
- `wgpu-types`'s `trace` and `replay` features have been replaced by the `serde` feature. By @KirmesBude in [#5149](https://github.com/gfx-rs/wgpu/pull/5149)
|
||||
|
||||
@ -895,6 +895,15 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
.expect("bit count overflowed 32 bits, somehow!?")])
|
||||
})
|
||||
}
|
||||
crate::MathFunction::CountTrailingZeros => {
|
||||
component_wise_concrete_int!(self, span, [arg], |e| {
|
||||
#[allow(clippy::useless_conversion)]
|
||||
Ok([e
|
||||
.trailing_zeros()
|
||||
.try_into()
|
||||
.expect("bit count overflowed 32 bits, somehow!?")])
|
||||
})
|
||||
}
|
||||
crate::MathFunction::Floor => {
|
||||
component_wise_float!(self, span, [arg], |e| { Ok([e.floor()]) })
|
||||
}
|
||||
|
||||
@ -75,14 +75,10 @@ void main() {
|
||||
uint ftb_b = uint(findLSB(1u));
|
||||
ivec2 ftb_c = findLSB(ivec2(-1));
|
||||
uvec2 ftb_d = uvec2(findLSB(uvec2(1u)));
|
||||
uint ctz_a = min(uint(findLSB(0u)), 32u);
|
||||
int ctz_b = int(min(uint(findLSB(0)), 32u));
|
||||
uint ctz_c = min(uint(findLSB(4294967295u)), 32u);
|
||||
int ctz_d = int(min(uint(findLSB(-1)), 32u));
|
||||
uvec2 ctz_e = min(uvec2(findLSB(uvec2(0u))), uvec2(32u));
|
||||
ivec2 ctz_f = ivec2(min(uvec2(findLSB(ivec2(0))), uvec2(32u)));
|
||||
uvec2 ctz_g = min(uvec2(findLSB(uvec2(1u))), uvec2(32u));
|
||||
ivec2 ctz_h = ivec2(min(uvec2(findLSB(ivec2(1))), uvec2(32u)));
|
||||
uvec2 ctz_e = uvec2(32u, 32u);
|
||||
ivec2 ctz_f = ivec2(32, 32);
|
||||
uvec2 ctz_g = uvec2(0u, 0u);
|
||||
ivec2 ctz_h = ivec2(0, 0);
|
||||
ivec2 clz_c = ivec2(0, 0);
|
||||
uvec2 clz_d = uvec2(31u, 31u);
|
||||
float lde_a = ldexp(1.0, 2);
|
||||
|
||||
@ -85,14 +85,10 @@ void main()
|
||||
uint ftb_b = firstbitlow(1u);
|
||||
int2 ftb_c = asint(firstbitlow((-1).xx));
|
||||
uint2 ftb_d = firstbitlow((1u).xx);
|
||||
uint ctz_a = min(32u, firstbitlow(0u));
|
||||
int ctz_b = asint(min(32u, firstbitlow(0)));
|
||||
uint ctz_c = min(32u, firstbitlow(4294967295u));
|
||||
int ctz_d = asint(min(32u, firstbitlow(-1)));
|
||||
uint2 ctz_e = min((32u).xx, firstbitlow((0u).xx));
|
||||
int2 ctz_f = asint(min((32u).xx, firstbitlow((0).xx)));
|
||||
uint2 ctz_g = min((32u).xx, firstbitlow((1u).xx));
|
||||
int2 ctz_h = asint(min((32u).xx, firstbitlow((1).xx)));
|
||||
uint2 ctz_e = uint2(32u, 32u);
|
||||
int2 ctz_f = int2(32, 32);
|
||||
uint2 ctz_g = uint2(0u, 0u);
|
||||
int2 ctz_h = int2(0, 0);
|
||||
int2 clz_c = int2(0, 0);
|
||||
uint2 clz_d = uint2(31u, 31u);
|
||||
float lde_a = ldexp(1.0, 2);
|
||||
|
||||
@ -80,14 +80,10 @@ fragment void main_(
|
||||
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);
|
||||
uint ctz_a = metal::ctz(0u);
|
||||
int ctz_b = metal::ctz(0);
|
||||
uint ctz_c = metal::ctz(4294967295u);
|
||||
int ctz_d = metal::ctz(-1);
|
||||
metal::uint2 ctz_e = metal::ctz(metal::uint2(0u));
|
||||
metal::int2 ctz_f = metal::ctz(metal::int2(0));
|
||||
metal::uint2 ctz_g = metal::ctz(metal::uint2(1u));
|
||||
metal::int2 ctz_h = metal::ctz(metal::int2(1));
|
||||
metal::uint2 ctz_e = metal::uint2(32u, 32u);
|
||||
metal::int2 ctz_f = metal::int2(32, 32);
|
||||
metal::uint2 ctz_g = metal::uint2(0u, 0u);
|
||||
metal::int2 ctz_h = metal::int2(0, 0);
|
||||
metal::int2 clz_c = metal::int2(0, 0);
|
||||
metal::uint2 clz_d = metal::uint2(31u, 31u);
|
||||
float lde_a = metal::ldexp(1.0, 2);
|
||||
|
||||
@ -1,7 +1,7 @@
|
||||
; SPIR-V
|
||||
; Version: 1.1
|
||||
; Generator: rspirv
|
||||
; Bound: 118
|
||||
; Bound: 100
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
@ -44,15 +44,15 @@ OpMemberDecorate %15 1 Offset 16
|
||||
%28 = OpConstantComposite %5 %22 %22
|
||||
%29 = OpConstant %8 1
|
||||
%30 = OpConstantComposite %7 %29 %29
|
||||
%31 = OpConstant %6 0
|
||||
%32 = OpConstant %8 4294967295
|
||||
%33 = OpConstantComposite %7 %27 %27
|
||||
%34 = OpConstantComposite %5 %31 %31
|
||||
%35 = OpConstant %6 1
|
||||
%36 = OpConstantComposite %5 %35 %35
|
||||
%37 = OpConstant %8 31
|
||||
%38 = OpConstantComposite %5 %31 %31
|
||||
%39 = OpConstantComposite %7 %37 %37
|
||||
%31 = OpConstant %8 32
|
||||
%32 = OpConstant %6 32
|
||||
%33 = OpConstant %6 0
|
||||
%34 = OpConstantComposite %7 %31 %31
|
||||
%35 = OpConstantComposite %5 %32 %32
|
||||
%36 = OpConstantComposite %7 %27 %27
|
||||
%37 = OpConstantComposite %5 %33 %33
|
||||
%38 = OpConstant %8 31
|
||||
%39 = OpConstantComposite %7 %38 %38
|
||||
%40 = OpConstant %6 2
|
||||
%41 = OpConstant %4 2.0
|
||||
%42 = OpConstantComposite %9 %19 %41
|
||||
@ -64,8 +64,6 @@ OpMemberDecorate %15 1 Offset 16
|
||||
%48 = OpConstantComposite %3 %46 %46 %46 %46
|
||||
%55 = OpConstantComposite %3 %19 %19 %19 %19
|
||||
%62 = OpConstantNull %6
|
||||
%79 = OpConstant %8 32
|
||||
%88 = OpConstantComposite %7 %79 %79
|
||||
%17 = OpFunction %2 None %18
|
||||
%16 = OpLabel
|
||||
OpBranch %49
|
||||
@ -96,43 +94,27 @@ OpBranch %49
|
||||
%75 = OpExtInst %8 %1 FindILsb %29
|
||||
%76 = OpExtInst %5 %1 FindILsb %28
|
||||
%77 = OpExtInst %7 %1 FindILsb %30
|
||||
%80 = OpExtInst %8 %1 FindILsb %27
|
||||
%78 = OpExtInst %8 %1 UMin %79 %80
|
||||
%82 = OpExtInst %6 %1 FindILsb %31
|
||||
%81 = OpExtInst %6 %1 UMin %79 %82
|
||||
%84 = OpExtInst %8 %1 FindILsb %32
|
||||
%83 = OpExtInst %8 %1 UMin %79 %84
|
||||
%86 = OpExtInst %6 %1 FindILsb %22
|
||||
%85 = OpExtInst %6 %1 UMin %79 %86
|
||||
%89 = OpExtInst %7 %1 FindILsb %33
|
||||
%87 = OpExtInst %7 %1 UMin %88 %89
|
||||
%91 = OpExtInst %5 %1 FindILsb %34
|
||||
%90 = OpExtInst %5 %1 UMin %88 %91
|
||||
%93 = OpExtInst %7 %1 FindILsb %30
|
||||
%92 = OpExtInst %7 %1 UMin %88 %93
|
||||
%95 = OpExtInst %5 %1 FindILsb %36
|
||||
%94 = OpExtInst %5 %1 UMin %88 %95
|
||||
%96 = OpExtInst %4 %1 Ldexp %19 %40
|
||||
%97 = OpExtInst %9 %1 Ldexp %42 %45
|
||||
%98 = OpExtInst %10 %1 ModfStruct %46
|
||||
%99 = OpExtInst %10 %1 ModfStruct %46
|
||||
%100 = OpCompositeExtract %4 %99 0
|
||||
%101 = OpExtInst %10 %1 ModfStruct %46
|
||||
%102 = OpCompositeExtract %4 %101 1
|
||||
%103 = OpExtInst %11 %1 ModfStruct %47
|
||||
%104 = OpExtInst %12 %1 ModfStruct %48
|
||||
%105 = OpCompositeExtract %3 %104 1
|
||||
%106 = OpCompositeExtract %4 %105 0
|
||||
%107 = OpExtInst %11 %1 ModfStruct %47
|
||||
%108 = OpCompositeExtract %9 %107 0
|
||||
%109 = OpCompositeExtract %4 %108 1
|
||||
%110 = OpExtInst %13 %1 FrexpStruct %46
|
||||
%111 = OpExtInst %13 %1 FrexpStruct %46
|
||||
%112 = OpCompositeExtract %4 %111 0
|
||||
%113 = OpExtInst %13 %1 FrexpStruct %46
|
||||
%114 = OpCompositeExtract %6 %113 1
|
||||
%115 = OpExtInst %15 %1 FrexpStruct %48
|
||||
%116 = OpCompositeExtract %14 %115 1
|
||||
%117 = OpCompositeExtract %6 %116 0
|
||||
%78 = OpExtInst %4 %1 Ldexp %19 %40
|
||||
%79 = OpExtInst %9 %1 Ldexp %42 %45
|
||||
%80 = OpExtInst %10 %1 ModfStruct %46
|
||||
%81 = OpExtInst %10 %1 ModfStruct %46
|
||||
%82 = OpCompositeExtract %4 %81 0
|
||||
%83 = OpExtInst %10 %1 ModfStruct %46
|
||||
%84 = OpCompositeExtract %4 %83 1
|
||||
%85 = OpExtInst %11 %1 ModfStruct %47
|
||||
%86 = OpExtInst %12 %1 ModfStruct %48
|
||||
%87 = OpCompositeExtract %3 %86 1
|
||||
%88 = OpCompositeExtract %4 %87 0
|
||||
%89 = OpExtInst %11 %1 ModfStruct %47
|
||||
%90 = OpCompositeExtract %9 %89 0
|
||||
%91 = OpCompositeExtract %4 %90 1
|
||||
%92 = OpExtInst %13 %1 FrexpStruct %46
|
||||
%93 = OpExtInst %13 %1 FrexpStruct %46
|
||||
%94 = OpCompositeExtract %4 %93 0
|
||||
%95 = OpExtInst %13 %1 FrexpStruct %46
|
||||
%96 = OpCompositeExtract %6 %95 1
|
||||
%97 = OpExtInst %15 %1 FrexpStruct %48
|
||||
%98 = OpCompositeExtract %14 %97 1
|
||||
%99 = OpCompositeExtract %6 %98 0
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -20,14 +20,10 @@ fn main() {
|
||||
let ftb_b = firstTrailingBit(1u);
|
||||
let ftb_c = firstTrailingBit(vec2(-1i));
|
||||
let ftb_d = firstTrailingBit(vec2(1u));
|
||||
let ctz_a = countTrailingZeros(0u);
|
||||
let ctz_b = countTrailingZeros(0i);
|
||||
let ctz_c = countTrailingZeros(4294967295u);
|
||||
let ctz_d = countTrailingZeros(-1i);
|
||||
let ctz_e = countTrailingZeros(vec2(0u));
|
||||
let ctz_f = countTrailingZeros(vec2(0i));
|
||||
let ctz_g = countTrailingZeros(vec2(1u));
|
||||
let ctz_h = countTrailingZeros(vec2(1i));
|
||||
let ctz_e = vec2<u32>(32u, 32u);
|
||||
let ctz_f = vec2<i32>(32i, 32i);
|
||||
let ctz_g = vec2<u32>(0u, 0u);
|
||||
let ctz_h = vec2<i32>(0i, 0i);
|
||||
let clz_c = vec2<i32>(0i, 0i);
|
||||
let clz_d = vec2<u32>(31u, 31u);
|
||||
let lde_a = ldexp(1f, 2i);
|
||||
|
||||
Loading…
Reference in New Issue
Block a user