mirror of
https://github.com/gfx-rs/wgpu.git
synced 2024-11-25 16:24:24 +00:00
feat(const_eval): impl. sign with new component_wise_signed
This commit is contained in:
Erich Gubler
Teodor Tanasoaia
parent
0bd5f77601
commit
2d3005b745
@ -82,6 +82,7 @@ Bottom level categories:
|
||||
- `log2`
|
||||
- `radians`
|
||||
- `reverseBits`
|
||||
- `sign`
|
||||
- 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)
|
||||
- `wgpu-core`'s `serial-pass` feature has been removed. Use `serde` instead. By @KirmesBude in [#5149](https://github.com/gfx-rs/wgpu/pull/5149)
|
||||
|
||||
@ -234,6 +234,22 @@ gen_component_wise_extractor! {
|
||||
],
|
||||
}
|
||||
|
||||
gen_component_wise_extractor! {
|
||||
component_wise_signed -> Signed,
|
||||
literals: [
|
||||
AbstractFloat => AbstractFloat: f64,
|
||||
AbstractInt => AbstractInt: i64,
|
||||
F32 => F32: f32,
|
||||
I32 => I32: i32,
|
||||
],
|
||||
scalar_kinds: [
|
||||
Sint,
|
||||
AbstractInt,
|
||||
Float,
|
||||
AbstractFloat,
|
||||
],
|
||||
}
|
||||
|
||||
#[derive(Debug)]
|
||||
enum Behavior {
|
||||
Wgsl,
|
||||
@ -975,6 +991,9 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
crate::MathFunction::Saturate => {
|
||||
component_wise_float!(self, span, [arg], |e| { Ok([e.clamp(0., 1.)]) })
|
||||
}
|
||||
crate::MathFunction::Sign => {
|
||||
component_wise_signed!(self, span, [arg], |e| { Ok([e.signum()]) })
|
||||
}
|
||||
crate::MathFunction::Sin => {
|
||||
component_wise_float!(self, span, [arg], |e| { Ok([e.sin()]) })
|
||||
}
|
||||
|
||||
@ -62,10 +62,8 @@ void main() {
|
||||
vec4 d = radians(v);
|
||||
vec4 e = clamp(v, vec4(0.0), vec4(1.0));
|
||||
vec4 g = refract(v, v, 1.0);
|
||||
int sign_a = sign(-1);
|
||||
ivec4 sign_b = sign(ivec4(-1));
|
||||
float sign_c = sign(-1.0);
|
||||
vec4 sign_d = sign(vec4(-1.0));
|
||||
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);
|
||||
|
||||
@ -72,10 +72,8 @@ void main()
|
||||
float4 d = radians(v);
|
||||
float4 e = saturate(v);
|
||||
float4 g = refract(v, v, 1.0);
|
||||
int sign_a = sign(-1);
|
||||
int4 sign_b = sign((-1).xxxx);
|
||||
float sign_c = sign(-1.0);
|
||||
float4 sign_d = sign((-1.0).xxxx);
|
||||
int4 sign_b = int4(-1, -1, -1, -1);
|
||||
float4 sign_d = float4(-1.0, -1.0, -1.0, -1.0);
|
||||
int const_dot = dot((int2)0, (int2)0);
|
||||
uint first_leading_bit_abs = firstbithigh(0u);
|
||||
int flb_a = asint(firstbithigh(-1));
|
||||
|
||||
@ -64,18 +64,15 @@ fragment void main_(
|
||||
metal::float4 d = ((v) * 0.017453292519943295474);
|
||||
metal::float4 e = metal::saturate(v);
|
||||
metal::float4 g = metal::refract(v, v, 1.0);
|
||||
int sign_a = metal::select(metal::select(-1, 1, (-1 > 0)), 0, (-1 == 0));
|
||||
metal::int4 _e12 = metal::int4(-1);
|
||||
metal::int4 sign_b = metal::select(metal::select(int4(-1), int4(1), (_e12 > 0)), 0, (_e12 == 0));
|
||||
float sign_c = metal::sign(-1.0);
|
||||
metal::float4 sign_d = metal::sign(metal::float4(-1.0));
|
||||
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 _e27 = metal::int2(-1);
|
||||
metal::int2 flb_b = metal::select(31 - metal::clz(metal::select(_e27, ~_e27, _e27 < 0)), int2(-1), _e27 == 0 || _e27 == -1);
|
||||
metal::uint2 _e30 = metal::uint2(1u);
|
||||
metal::uint2 flb_c = metal::select(31 - metal::clz(_e30), uint2(-1), _e30 == 0 || _e30 == -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);
|
||||
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);
|
||||
|
||||
@ -1,69 +1,69 @@
|
||||
; SPIR-V
|
||||
; Version: 1.1
|
||||
; Generator: rspirv
|
||||
; Bound: 100
|
||||
; Bound: 96
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint Fragment %17 "main"
|
||||
OpExecutionMode %17 OriginUpperLeft
|
||||
OpMemberDecorate %10 0 Offset 0
|
||||
OpMemberDecorate %10 1 Offset 4
|
||||
OpMemberDecorate %11 0 Offset 0
|
||||
OpMemberDecorate %11 1 Offset 8
|
||||
OpMemberDecorate %11 1 Offset 4
|
||||
OpMemberDecorate %12 0 Offset 0
|
||||
OpMemberDecorate %12 1 Offset 16
|
||||
OpMemberDecorate %12 1 Offset 8
|
||||
OpMemberDecorate %13 0 Offset 0
|
||||
OpMemberDecorate %13 1 Offset 4
|
||||
OpMemberDecorate %13 1 Offset 16
|
||||
OpMemberDecorate %14 0 Offset 0
|
||||
OpMemberDecorate %14 1 Offset 4
|
||||
OpMemberDecorate %15 0 Offset 0
|
||||
OpMemberDecorate %15 1 Offset 16
|
||||
%2 = OpTypeVoid
|
||||
%4 = OpTypeFloat 32
|
||||
%3 = OpTypeVector %4 4
|
||||
%6 = OpTypeInt 32 1
|
||||
%5 = OpTypeVector %6 2
|
||||
%8 = OpTypeInt 32 0
|
||||
%7 = OpTypeVector %8 2
|
||||
%9 = OpTypeVector %4 2
|
||||
%10 = OpTypeStruct %4 %4
|
||||
%11 = OpTypeStruct %9 %9
|
||||
%12 = OpTypeStruct %3 %3
|
||||
%13 = OpTypeStruct %4 %6
|
||||
%14 = OpTypeVector %6 4
|
||||
%15 = OpTypeStruct %3 %14
|
||||
%5 = OpTypeVector %6 4
|
||||
%7 = OpTypeVector %6 2
|
||||
%9 = OpTypeInt 32 0
|
||||
%8 = OpTypeVector %9 2
|
||||
%10 = OpTypeVector %4 2
|
||||
%11 = OpTypeStruct %4 %4
|
||||
%12 = OpTypeStruct %10 %10
|
||||
%13 = OpTypeStruct %3 %3
|
||||
%14 = OpTypeStruct %4 %6
|
||||
%15 = OpTypeStruct %3 %5
|
||||
%18 = OpTypeFunction %2
|
||||
%19 = OpConstant %4 1.0
|
||||
%20 = OpConstant %4 0.0
|
||||
%21 = OpConstantComposite %3 %20 %20 %20 %20
|
||||
%22 = OpConstant %6 -1
|
||||
%23 = OpConstantComposite %14 %22 %22 %22 %22
|
||||
%23 = OpConstantComposite %5 %22 %22 %22 %22
|
||||
%24 = OpConstant %4 -1.0
|
||||
%25 = OpConstantComposite %3 %24 %24 %24 %24
|
||||
%26 = OpConstantNull %5
|
||||
%27 = OpConstant %8 0
|
||||
%28 = OpConstantComposite %5 %22 %22
|
||||
%29 = OpConstant %8 1
|
||||
%30 = OpConstantComposite %7 %29 %29
|
||||
%31 = OpConstant %8 32
|
||||
%26 = OpConstantNull %7
|
||||
%27 = OpConstant %9 0
|
||||
%28 = OpConstantComposite %7 %22 %22
|
||||
%29 = OpConstant %9 1
|
||||
%30 = OpConstantComposite %8 %29 %29
|
||||
%31 = OpConstant %9 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
|
||||
%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 %9 %19 %41
|
||||
%42 = OpConstantComposite %10 %19 %41
|
||||
%43 = OpConstant %6 3
|
||||
%44 = OpConstant %6 4
|
||||
%45 = OpConstantComposite %5 %43 %44
|
||||
%45 = OpConstantComposite %7 %43 %44
|
||||
%46 = OpConstant %4 1.5
|
||||
%47 = OpConstantComposite %9 %46 %46
|
||||
%47 = OpConstantComposite %10 %46 %46
|
||||
%48 = OpConstantComposite %3 %46 %46 %46 %46
|
||||
%55 = OpConstantComposite %3 %19 %19 %19 %19
|
||||
%62 = OpConstantNull %6
|
||||
%58 = OpConstantNull %6
|
||||
%17 = OpFunction %2 None %18
|
||||
%16 = OpLabel
|
||||
OpBranch %49
|
||||
@ -74,47 +74,43 @@ OpBranch %49
|
||||
%53 = OpExtInst %3 %1 Radians %21
|
||||
%54 = OpExtInst %3 %1 FClamp %21 %21 %55
|
||||
%56 = OpExtInst %3 %1 Refract %21 %21 %19
|
||||
%57 = OpExtInst %6 %1 SSign %22
|
||||
%58 = OpExtInst %14 %1 SSign %23
|
||||
%59 = OpExtInst %4 %1 FSign %24
|
||||
%60 = OpExtInst %3 %1 FSign %25
|
||||
%63 = OpCompositeExtract %6 %26 0
|
||||
%64 = OpCompositeExtract %6 %26 0
|
||||
%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
|
||||
%66 = OpIAdd %6 %62 %65
|
||||
%67 = OpCompositeExtract %6 %26 1
|
||||
%68 = OpCompositeExtract %6 %26 1
|
||||
%69 = OpIMul %6 %67 %68
|
||||
%61 = OpIAdd %6 %66 %69
|
||||
%70 = OpExtInst %8 %1 FindUMsb %27
|
||||
%71 = OpExtInst %6 %1 FindSMsb %22
|
||||
%72 = OpExtInst %5 %1 FindSMsb %28
|
||||
%73 = OpExtInst %7 %1 FindUMsb %30
|
||||
%74 = OpExtInst %6 %1 FindILsb %22
|
||||
%75 = OpExtInst %8 %1 FindILsb %29
|
||||
%76 = OpExtInst %5 %1 FindILsb %28
|
||||
%77 = OpExtInst %7 %1 FindILsb %30
|
||||
%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
|
||||
%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
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -7,10 +7,8 @@ fn main() {
|
||||
let d = radians(v);
|
||||
let e = saturate(v);
|
||||
let g = refract(v, v, 1f);
|
||||
let sign_a = sign(-1i);
|
||||
let sign_b = sign(vec4(-1i));
|
||||
let sign_c = sign(-1f);
|
||||
let sign_d = sign(vec4(-1f));
|
||||
let sign_b = vec4<i32>(-1i, -1i, -1i, -1i);
|
||||
let sign_d = vec4<f32>(-1f, -1f, -1f, -1f);
|
||||
let const_dot = dot(vec2<i32>(), vec2<i32>());
|
||||
let first_leading_bit_abs = firstLeadingBit(0u);
|
||||
let flb_a = firstLeadingBit(-1i);
|
||||
|
||||
Loading…
Reference in New Issue
Block a user