From 974093503e8ac46efdbd1cf9ec589b9176c4650b Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Sat, 28 Jan 2023 11:18:37 -0500 Subject: [PATCH 01/10] Add countLeadingZeros --- src/back/glsl/mod.rs | 34 ++++++++ src/back/hlsl/writer.rs | 39 ++++++++++ src/back/msl/writer.rs | 1 + src/back/spv/block.rs | 69 +++++++++++++++- src/back/wgsl/writer.rs | 1 + src/front/wgsl/parse/conv.rs | 1 + src/lib.rs | 1 + src/proc/mod.rs | 1 + src/proc/typifier.rs | 1 + src/valid/expression.rs | 6 +- tests/in/math-functions.wgsl | 4 + .../out/glsl/math-functions.main.Vertex.glsl | 4 + tests/out/hlsl/math-functions.hlsl | 4 + tests/out/msl/math-functions.msl | 4 + tests/out/spv/math-functions.spvasm | 78 ++++++++++++------- tests/out/wgsl/math-functions.wgsl | 4 + 16 files changed, 219 insertions(+), 33 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 4208856387..cc4e851d1b 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -2928,6 +2928,40 @@ impl<'a, W: Write> Writer<'a, W> { Mf::Transpose => "transpose", Mf::Determinant => "determinant", // bits + Mf::CountLeadingZeros => { + match *ctx.info[arg].ty.inner_with(&self.module.types) { + crate::TypeInner::Vector { size, kind, .. } => { + let s = back::vector_size_str(size); + + if let crate::ScalarKind::Uint = kind { + write!(self.out, "uvec{s}(ivec{s}(31) - findMSB(")?; + self.write_expr(arg, ctx)?; + write!(self.out, "))")?; + } else { + write!(self.out, "mix(ivec{s}(31) - findMSB(")?; + self.write_expr(arg, ctx)?; + write!(self.out, "), ivec{s}(0), lessThan(")?; + self.write_expr(arg, ctx)?; + write!(self.out, ", ivec{s}(0)))")?; + } + } + crate::TypeInner::Scalar { kind, .. } => { + if let crate::ScalarKind::Uint = kind { + write!(self.out, "uint(31 - findMSB(")?; + } else { + write!(self.out, "(")?; + self.write_expr(arg, ctx)?; + write!(self.out, " < 0 ? 0 : 31 - findMSB(")?; + } + + self.write_expr(arg, ctx)?; + write!(self.out, "))")?; + } + _ => unreachable!(), + }; + + return Ok(()); + } Mf::CountOneBits => "bitCount", Mf::ReverseBits => "bitfieldReverse", Mf::ExtractBits => "bitfieldExtract", diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 7ac5563aaa..62136a61db 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -2510,6 +2510,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { Unpack2x16float, Regular(&'static str), MissingIntOverload(&'static str), + CountLeadingZeros, } let fun = match fun { @@ -2572,6 +2573,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { Mf::Transpose => Function::Regular("transpose"), Mf::Determinant => Function::Regular("determinant"), // bits + Mf::CountLeadingZeros => Function::CountLeadingZeros, Mf::CountOneBits => Function::MissingIntOverload("countbits"), Mf::ReverseBits => Function::MissingIntOverload("reversebits"), Mf::FindLsb => Function::Regular("firstbitlow"), @@ -2639,6 +2641,43 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { write!(self.out, ")")?; } } + Function::CountLeadingZeros => { + match *func_ctx.info[arg].ty.inner_with(&module.types) { + TypeInner::Vector { size, kind, .. } => { + let s = match size { + crate::VectorSize::Bi => ".xx", + crate::VectorSize::Tri => ".xxx", + crate::VectorSize::Quad => ".xxxx", + }; + + if let ScalarKind::Uint = kind { + write!(self.out, "asuint((31){s} - firstbithigh(")?; + } else { + write!(self.out, "(")?; + self.write_expr(module, arg, func_ctx)?; + write!( + self.out, + " < (0){s} ? (0){s} : (31){s} - firstbithigh(" + )?; + } + } + TypeInner::Scalar { kind, .. } => { + if let ScalarKind::Uint = kind { + write!(self.out, "asuint(31 - firstbithigh(")?; + } else { + write!(self.out, "(")?; + self.write_expr(module, arg, func_ctx)?; + write!(self.out, " < 0 ? 0 : 31 - firstbithigh(")?; + } + } + _ => unreachable!(), + } + + self.write_expr(module, arg, func_ctx)?; + write!(self.out, "))")?; + + return Ok(()); + } } } Expression::Swizzle { diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index b72c76783b..04c1fed213 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -1689,6 +1689,7 @@ impl Writer { Mf::Transpose => "transpose", Mf::Determinant => "determinant", // bits + Mf::CountLeadingZeros => "clz", Mf::CountOneBits => "popcount", Mf::ReverseBits => "reverse_bits", Mf::ExtractBits => "extract_bits", diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index 5ebbf2ea88..bf09fc3ad1 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -724,7 +724,7 @@ impl<'w> BlockContext<'w> { self.temp_list.resize(size as _, arg1_id); let id = self.gen_id(); - block.body.push(Instruction::composite_construct( + block.body.push(Instruction::constant_composite( result_type_id, id, &self.temp_list, @@ -735,7 +735,7 @@ impl<'w> BlockContext<'w> { self.temp_list.resize(size as _, arg2_id); let id = self.gen_id(); - block.body.push(Instruction::composite_construct( + block.body.push(Instruction::constant_composite( result_type_id, id, &self.temp_list, @@ -888,6 +888,71 @@ impl<'w> BlockContext<'w> { id, arg0_id, )), + Mf::CountLeadingZeros => { + let int = crate::ScalarValue::Sint(31); + + let (int_type_id, int_id) = match *arg_ty { + crate::TypeInner::Vector { size, width, .. } => { + let ty = self.get_type_id(LookupType::Local(LocalType::Value { + vector_size: Some(size), + kind: crate::ScalarKind::Sint, + width, + pointer_space: None, + })); + + self.temp_list.clear(); + self.temp_list + .resize(size as _, self.writer.get_constant_scalar(int, width)); + + let id = self.gen_id(); + block.body.push(Instruction::constant_composite( + ty, + id, + &self.temp_list, + )); + + (ty, id) + } + crate::TypeInner::Scalar { width, .. } => ( + self.get_type_id(LookupType::Local(LocalType::Value { + vector_size: None, + kind: crate::ScalarKind::Sint, + width, + pointer_space: None, + })), + self.writer.get_constant_scalar(int, width), + ), + _ => unreachable!(), + }; + + block.body.push(Instruction::ext_inst( + self.writer.gl450_ext_inst_id, + spirv::GLOp::FindUMsb, + int_type_id, + id, + &[arg0_id], + )); + + let sub_id = self.gen_id(); + block.body.push(Instruction::binary( + spirv::Op::ISub, + int_type_id, + sub_id, + int_id, + id, + )); + + if let Some(crate::ScalarKind::Uint) = arg_scalar_kind { + block.body.push(Instruction::unary( + spirv::Op::Bitcast, + result_type_id, + self.gen_id(), + sub_id, + )); + } + + return Ok(()); + } Mf::CountOneBits => MathOp::Custom(Instruction::unary( spirv::Op::BitCount, result_type_id, diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index 5524ca9c20..4f58cd9ee8 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -1578,6 +1578,7 @@ impl Writer { Mf::Transpose => Function::Regular("transpose"), Mf::Determinant => Function::Regular("determinant"), // bits + Mf::CountLeadingZeros => Function::Regular("countLeadingZeros"), Mf::CountOneBits => Function::Regular("countOneBits"), Mf::ReverseBits => Function::Regular("reverseBits"), Mf::ExtractBits => Function::Regular("extractBits"), diff --git a/src/front/wgsl/parse/conv.rs b/src/front/wgsl/parse/conv.rs index ca4cc5519e..4164332166 100644 --- a/src/front/wgsl/parse/conv.rs +++ b/src/front/wgsl/parse/conv.rs @@ -191,6 +191,7 @@ pub fn map_standard_fun(word: &str) -> Option { "transpose" => Mf::Transpose, "determinant" => Mf::Determinant, // bits + "countLeadingZeros" => Mf::CountLeadingZeros, "countOneBits" => Mf::CountOneBits, "reverseBits" => Mf::ReverseBits, "extractBits" => Mf::ExtractBits, diff --git a/src/lib.rs b/src/lib.rs index ceae0c14ca..136205ca17 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -1066,6 +1066,7 @@ pub enum MathFunction { Transpose, Determinant, // bits + CountLeadingZeros, CountOneBits, ReverseBits, ExtractBits, diff --git a/src/proc/mod.rs b/src/proc/mod.rs index f30158c113..0a8e4a961a 100644 --- a/src/proc/mod.rs +++ b/src/proc/mod.rs @@ -279,6 +279,7 @@ impl super::MathFunction { Self::Transpose => 1, Self::Determinant => 1, // bits + Self::CountLeadingZeros => 1, Self::CountOneBits => 1, Self::ReverseBits => 1, Self::ExtractBits => 3, diff --git a/src/proc/typifier.rs b/src/proc/typifier.rs index 4c55d101ab..64896ec413 100644 --- a/src/proc/typifier.rs +++ b/src/proc/typifier.rs @@ -793,6 +793,7 @@ impl<'a> ResolveContext<'a> { )), }, // bits + Mf::CountLeadingZeros | Mf::CountOneBits | Mf::ReverseBits | Mf::ExtractBits | diff --git a/src/valid/expression.rs b/src/valid/expression.rs index 9853f1d1be..d599e6b62f 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -1223,7 +1223,11 @@ impl super::Validator { _ => return Err(ExpressionError::InvalidArgumentType(fun, 0, arg)), } } - Mf::CountOneBits | Mf::ReverseBits | Mf::FindLsb | Mf::FindMsb => { + Mf::CountLeadingZeros + | Mf::CountOneBits + | Mf::ReverseBits + | Mf::FindLsb + | Mf::FindMsb => { if arg1_ty.is_some() | arg2_ty.is_some() | arg3_ty.is_some() { return Err(ExpressionError::WrongArgumentCount(fun)); } diff --git a/tests/in/math-functions.wgsl b/tests/in/math-functions.wgsl index 13b080e927..1c2a8e4579 100644 --- a/tests/in/math-functions.wgsl +++ b/tests/in/math-functions.wgsl @@ -10,4 +10,8 @@ fn main() { let g = refract(v, v, f); let const_dot = dot(vec2(), vec2()); let first_leading_bit_abs = firstLeadingBit(abs(0u)); + let clz_a = countLeadingZeros(-1); + let clz_b = countLeadingZeros(1u); + let clz_c = countLeadingZeros(vec2(-1)); + let clz_d = countLeadingZeros(vec2(1u)); } diff --git a/tests/out/glsl/math-functions.main.Vertex.glsl b/tests/out/glsl/math-functions.main.Vertex.glsl index 5efa8f27ca..46369f0930 100644 --- a/tests/out/glsl/math-functions.main.Vertex.glsl +++ b/tests/out/glsl/math-functions.main.Vertex.glsl @@ -14,5 +14,9 @@ void main() { vec4 g = refract(v, v, 1.0); int const_dot = ( + ivec2(0, 0).x * ivec2(0, 0).x + ivec2(0, 0).y * ivec2(0, 0).y); uint first_leading_bit_abs = uint(findMSB(uint(abs(int(0u))))); + int clz_a = (-1 < 0 ? 0 : 31 - findMSB(-1)); + uint clz_b = uint(31 - findMSB(1u)); + ivec2 clz_c = mix(ivec2(31) - findMSB(ivec2(-1)), ivec2(0), lessThan(ivec2(-1), ivec2(0))); + uvec2 clz_d = uvec2(ivec2(31) - findMSB(uvec2(1u))); } diff --git a/tests/out/hlsl/math-functions.hlsl b/tests/out/hlsl/math-functions.hlsl index 64c2e87eab..083da8ac8d 100644 --- a/tests/out/hlsl/math-functions.hlsl +++ b/tests/out/hlsl/math-functions.hlsl @@ -10,4 +10,8 @@ void main() float4 g = refract(v, v, 1.0); int const_dot = dot(int2(0, 0), int2(0, 0)); uint first_leading_bit_abs = firstbithigh(abs(0u)); + int clz_a = (-1 < 0 ? 0 : 31 - firstbithigh(-1)); + uint clz_b = asuint(31 - firstbithigh(1u)); + int2 clz_c = ((-1).xx < (0).xx ? (0).xx : (31).xx - firstbithigh((-1).xx)); + uint2 clz_d = asuint((31).xx - firstbithigh((1u).xx)); } diff --git a/tests/out/msl/math-functions.msl b/tests/out/msl/math-functions.msl index d097be47a8..0838a338c5 100644 --- a/tests/out/msl/math-functions.msl +++ b/tests/out/msl/math-functions.msl @@ -17,4 +17,8 @@ vertex void main_( metal::float4 g = metal::refract(v, v, 1.0); int const_dot = ( + const_type_1_.x * const_type_1_.x + const_type_1_.y * const_type_1_.y); uint first_leading_bit_abs = (((metal::clz(metal::abs(0u)) + 1) % 33) - 1); + int clz_a = metal::clz(-1); + uint clz_b = metal::clz(1u); + metal::int2 clz_c = metal::clz(metal::int2(-1)); + metal::uint2 clz_d = metal::clz(metal::uint2(1u)); } diff --git a/tests/out/spv/math-functions.spvasm b/tests/out/spv/math-functions.spvasm index 4fada56c7f..131c6caceb 100644 --- a/tests/out/spv/math-functions.spvasm +++ b/tests/out/spv/math-functions.spvasm @@ -1,11 +1,11 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 37 +; Bound: 55 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %14 "main" +OpEntryPoint Vertex %16 "main" %2 = OpTypeVoid %4 = OpTypeFloat 32 %3 = OpConstant %4 1.0 @@ -14,33 +14,51 @@ OpEntryPoint Vertex %14 "main" %6 = OpConstant %7 0 %9 = OpTypeInt 32 0 %8 = OpConstant %9 0 -%10 = OpTypeVector %4 4 -%11 = OpTypeVector %7 2 -%12 = OpConstantComposite %11 %6 %6 -%15 = OpTypeFunction %2 -%27 = OpConstantNull %7 -%14 = OpFunction %2 None %15 -%13 = OpLabel -OpBranch %16 -%16 = OpLabel -%17 = OpCompositeConstruct %10 %5 %5 %5 %5 -%18 = OpExtInst %4 %1 Degrees %3 -%19 = OpExtInst %4 %1 Radians %3 -%20 = OpExtInst %10 %1 Degrees %17 -%21 = OpExtInst %10 %1 Radians %17 -%23 = OpCompositeConstruct %10 %5 %5 %5 %5 -%24 = OpCompositeConstruct %10 %3 %3 %3 %3 -%22 = OpExtInst %10 %1 FClamp %17 %23 %24 -%25 = OpExtInst %10 %1 Refract %17 %17 %3 -%28 = OpCompositeExtract %7 %12 0 -%29 = OpCompositeExtract %7 %12 0 -%30 = OpIMul %7 %28 %29 -%31 = OpIAdd %7 %27 %30 -%32 = OpCompositeExtract %7 %12 1 -%33 = OpCompositeExtract %7 %12 1 -%34 = OpIMul %7 %32 %33 -%26 = OpIAdd %7 %31 %34 -%35 = OpCopyObject %9 %8 -%36 = OpExtInst %9 %1 FindUMsb %35 +%10 = OpConstant %7 -1 +%11 = OpConstant %9 1 +%12 = OpTypeVector %4 4 +%13 = OpTypeVector %7 2 +%14 = OpConstantComposite %13 %6 %6 +%17 = OpTypeFunction %2 +%40 = OpConstant %7 31 +%49 = OpTypeVector %9 2 +%25 = OpConstantComposite %12 %5 %5 %5 %5 +%26 = OpConstantComposite %12 %3 %3 %3 %3 +%29 = OpConstantNull %7 +%47 = OpConstantComposite %13 %40 %40 +%52 = OpConstantComposite %13 %40 %40 +%16 = OpFunction %2 None %17 +%15 = OpLabel +OpBranch %18 +%18 = OpLabel +%19 = OpCompositeConstruct %12 %5 %5 %5 %5 +%20 = OpExtInst %4 %1 Degrees %3 +%21 = OpExtInst %4 %1 Radians %3 +%22 = OpExtInst %12 %1 Degrees %19 +%23 = OpExtInst %12 %1 Radians %19 +%24 = OpExtInst %12 %1 FClamp %19 %25 %26 +%27 = OpExtInst %12 %1 Refract %19 %19 %3 +%30 = OpCompositeExtract %7 %14 0 +%31 = OpCompositeExtract %7 %14 0 +%32 = OpIMul %7 %30 %31 +%33 = OpIAdd %7 %29 %32 +%34 = OpCompositeExtract %7 %14 1 +%35 = OpCompositeExtract %7 %14 1 +%36 = OpIMul %7 %34 %35 +%28 = OpIAdd %7 %33 %36 +%37 = OpCopyObject %9 %8 +%38 = OpExtInst %9 %1 FindUMsb %37 +%39 = OpExtInst %7 %1 FindUMsb %10 +%41 = OpISub %7 %40 %39 +%42 = OpExtInst %7 %1 FindUMsb %11 +%43 = OpISub %7 %40 %42 +%44 = OpBitcast %9 %43 +%45 = OpCompositeConstruct %13 %10 %10 +%46 = OpExtInst %13 %1 FindUMsb %45 +%48 = OpISub %13 %47 %46 +%50 = OpCompositeConstruct %49 %11 %11 +%51 = OpExtInst %13 %1 FindUMsb %50 +%53 = OpISub %13 %52 %51 +%54 = OpBitcast %49 %53 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/math-functions.wgsl b/tests/out/wgsl/math-functions.wgsl index 004d9cf5be..d91a26cff4 100644 --- a/tests/out/wgsl/math-functions.wgsl +++ b/tests/out/wgsl/math-functions.wgsl @@ -9,4 +9,8 @@ fn main() { let g = refract(v, v, 1.0); let const_dot = dot(vec2(0, 0), vec2(0, 0)); let first_leading_bit_abs = firstLeadingBit(abs(0u)); + let clz_a = countLeadingZeros(-1); + let clz_b = countLeadingZeros(1u); + let clz_c = countLeadingZeros(vec2(-1)); + let clz_d = countLeadingZeros(vec2(1u)); } From e96edf5935555ee73a386b92a87d704929189034 Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Mon, 30 Jan 2023 15:23:01 -0500 Subject: [PATCH 02/10] [glsl-out] Bake countLeadingZeros --- src/back/glsl/mod.rs | 38 +++++++++++-------- .../out/glsl/math-functions.main.Vertex.glsl | 6 ++- 2 files changed, 27 insertions(+), 17 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index cc4e851d1b..704beca57c 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -1120,28 +1120,36 @@ impl<'a, W: Write> Writer<'a, W> { if min_ref_count <= expr_info.ref_count { self.need_bake_expressions.insert(expr.0); } - // if the expression is a Dot product with integer arguments, - // then the args needs baking as well - if let ( - fun_handle, - &Expression::Math { + + let (fun_handle, expr) = expr; + match *expr { + Expression::Math { fun: crate::MathFunction::Dot, arg, arg1, .. - }, - ) = expr - { - let inner = info[fun_handle].ty.inner_with(&self.module.types); - if let TypeInner::Scalar { kind, .. } = *inner { - match kind { - crate::ScalarKind::Sint | crate::ScalarKind::Uint => { - self.need_bake_expressions.insert(arg); - self.need_bake_expressions.insert(arg1.unwrap()); + } => { + // if the expression is a Dot product with integer arguments, + // then the args needs baking as well + let inner = info[fun_handle].ty.inner_with(&self.module.types); + if let TypeInner::Scalar { kind, .. } = *inner { + match kind { + crate::ScalarKind::Sint | crate::ScalarKind::Uint => { + self.need_bake_expressions.insert(arg); + self.need_bake_expressions.insert(arg1.unwrap()); + } + _ => (), } - _ => {} } } + Expression::Math { + fun: crate::MathFunction::CountLeadingZeros, + arg, + .. + } => { + self.need_bake_expressions.insert(arg); + } + _ => (), } } } diff --git a/tests/out/glsl/math-functions.main.Vertex.glsl b/tests/out/glsl/math-functions.main.Vertex.glsl index 46369f0930..b46e68de88 100644 --- a/tests/out/glsl/math-functions.main.Vertex.glsl +++ b/tests/out/glsl/math-functions.main.Vertex.glsl @@ -16,7 +16,9 @@ void main() { uint first_leading_bit_abs = uint(findMSB(uint(abs(int(0u))))); int clz_a = (-1 < 0 ? 0 : 31 - findMSB(-1)); uint clz_b = uint(31 - findMSB(1u)); - ivec2 clz_c = mix(ivec2(31) - findMSB(ivec2(-1)), ivec2(0), lessThan(ivec2(-1), ivec2(0))); - uvec2 clz_d = uvec2(ivec2(31) - findMSB(uvec2(1u))); + ivec2 _e20 = ivec2(-1); + ivec2 clz_c = mix(ivec2(31) - findMSB(_e20), ivec2(0), lessThan(_e20, ivec2(0))); + uvec2 _e23 = uvec2(1u); + uvec2 clz_d = uvec2(ivec2(31) - findMSB(_e23)); } From 727c06e452472549e509361d29c5dd6c17f54012 Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Mon, 30 Jan 2023 17:08:30 -0500 Subject: [PATCH 03/10] [hlsl-out] Bake countLeadingZeros --- src/back/glsl/mod.rs | 9 +++--- src/back/hlsl/mod.rs | 3 +- src/back/hlsl/writer.rs | 45 ++++++++++++++++++++++++------ tests/out/hlsl/math-functions.hlsl | 6 ++-- 4 files changed, 47 insertions(+), 16 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 704beca57c..eed219057b 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -1114,14 +1114,13 @@ impl<'a, W: Write> Writer<'a, W> { fn update_expressions_to_bake(&mut self, func: &crate::Function, info: &valid::FunctionInfo) { use crate::Expression; self.need_bake_expressions.clear(); - for expr in func.expressions.iter() { - let expr_info = &info[expr.0]; - let min_ref_count = func.expressions[expr.0].bake_ref_count(); + for (fun_handle, expr) in func.expressions.iter() { + let expr_info = &info[fun_handle]; + let min_ref_count = func.expressions[fun_handle].bake_ref_count(); if min_ref_count <= expr_info.ref_count { - self.need_bake_expressions.insert(expr.0); + self.need_bake_expressions.insert(fun_handle); } - let (fun_handle, expr) = expr; match *expr { Expression::Math { fun: crate::MathFunction::Dot, diff --git a/src/back/hlsl/mod.rs b/src/back/hlsl/mod.rs index 1031270390..6f31ed3743 100644 --- a/src/back/hlsl/mod.rs +++ b/src/back/hlsl/mod.rs @@ -107,7 +107,7 @@ mod writer; use std::fmt::Error as FmtError; use thiserror::Error; -use crate::proc; +use crate::{back, proc}; #[derive(Clone, Debug, Default, PartialEq, Eq, Hash)] #[cfg_attr(feature = "serialize", derive(serde::Serialize))] @@ -280,4 +280,5 @@ pub struct Writer<'a, W> { named_expressions: crate::NamedExpressions, wrapped: Wrapped, temp_access_chain: Vec, + need_bake_expressions: back::NeedBakeExpressions, } diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 62136a61db..1b314d41c5 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -83,6 +83,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { named_expressions: crate::NamedExpressions::default(), wrapped: super::Wrapped::default(), temp_access_chain: Vec::new(), + need_bake_expressions: Default::default(), } } @@ -93,6 +94,34 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { self.entry_point_io.clear(); self.named_expressions.clear(); self.wrapped.clear(); + self.need_bake_expressions.clear(); + } + + /// Helper method used to find which expressions of a given function require baking + /// + /// # Notes + /// Clears `need_bake_expressions` set before adding to it + fn update_expressions_to_bake(&mut self, func: &crate::Function, info: &valid::FunctionInfo) { + use crate::Expression; + self.need_bake_expressions.clear(); + for (fun_handle, expr) in func.expressions.iter() { + let expr_info = &info[fun_handle]; + let min_ref_count = func.expressions[fun_handle].bake_ref_count(); + if min_ref_count <= expr_info.ref_count { + self.need_bake_expressions.insert(fun_handle); + } + + match *expr { + Expression::Math { + fun: crate::MathFunction::CountLeadingZeros, + arg, + .. + } => { + self.need_bake_expressions.insert(arg); + } + _ => (), + } + } } pub fn write( @@ -244,7 +273,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { // before writing all statements and expressions. self.write_wrapped_functions(module, &ctx)?; - self.write_function(module, name.as_str(), function, &ctx)?; + self.write_function(module, name.as_str(), function, &ctx, info)?; writeln!(self.out)?; } @@ -296,7 +325,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } let name = self.names[&NameKey::EntryPoint(index as u16)].clone(); - self.write_function(module, &name, &ep.function, &ctx)?; + self.write_function(module, &name, &ep.function, &ctx, info)?; if index < module.entry_points.len() - 1 { writeln!(self.out)?; @@ -1034,9 +1063,12 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { name: &str, func: &crate::Function, func_ctx: &back::FunctionCtx<'_>, + info: &valid::FunctionInfo, ) -> BackendResult { // Function Declaration Syntax - https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-function-syntax + self.update_expressions_to_bake(func, info); + // Write modifier if let Some(crate::FunctionResult { binding: @@ -1284,15 +1316,12 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { // Otherwise, we could accidentally write variable name instead of full expression. // Also, we use sanitized names! It defense backend from generating variable with name from reserved keywords. Some(self.namer.call(name)) + } else if self.need_bake_expressions.contains(&handle) { + Some(format!("_expr{}", handle.index())) } else if info.ref_count == 0 { Some(self.namer.call("")) } else { - let min_ref_count = func_ctx.expressions[handle].bake_ref_count(); - if min_ref_count <= info.ref_count { - Some(format!("_expr{}", handle.index())) - } else { - None - } + None }; if let Some(name) = expr_name { diff --git a/tests/out/hlsl/math-functions.hlsl b/tests/out/hlsl/math-functions.hlsl index 083da8ac8d..0b97105598 100644 --- a/tests/out/hlsl/math-functions.hlsl +++ b/tests/out/hlsl/math-functions.hlsl @@ -12,6 +12,8 @@ void main() uint first_leading_bit_abs = firstbithigh(abs(0u)); int clz_a = (-1 < 0 ? 0 : 31 - firstbithigh(-1)); uint clz_b = asuint(31 - firstbithigh(1u)); - int2 clz_c = ((-1).xx < (0).xx ? (0).xx : (31).xx - firstbithigh((-1).xx)); - uint2 clz_d = asuint((31).xx - firstbithigh((1u).xx)); + int2 _expr20 = (-1).xx; + int2 clz_c = (_expr20 < (0).xx ? (0).xx : (31).xx - firstbithigh(_expr20)); + uint2 _expr23 = (1u).xx; + uint2 clz_d = asuint((31).xx - firstbithigh(_expr23)); } From 4a50966fe475f6c9533176a72b8bd12b37abb5e4 Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Mon, 30 Jan 2023 18:42:30 -0500 Subject: [PATCH 04/10] [hlsl-out] Update Baked expressions --- src/back/hlsl/writer.rs | 17 ++++++++++------- 1 file changed, 10 insertions(+), 7 deletions(-) diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 1b314d41c5..fd9b535796 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -112,13 +112,16 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } match *expr { - Expression::Math { - fun: crate::MathFunction::CountLeadingZeros, - arg, - .. - } => { - self.need_bake_expressions.insert(arg); - } + Expression::Math { fun, arg, .. } => match fun { + crate::MathFunction::Asinh + | crate::MathFunction::Acosh + | crate::MathFunction::Atanh + | crate::MathFunction::Unpack2x16float + | crate::MathFunction::CountLeadingZeros => { + self.need_bake_expressions.insert(arg); + } + _ => (), + }, _ => (), } } From d96521031f064cdca166772631f86e0baaed71e8 Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Mon, 30 Jan 2023 19:01:36 -0500 Subject: [PATCH 05/10] Remove unnecessary bake for sints --- src/back/glsl/mod.rs | 5 ++++- src/back/hlsl/writer.rs | 18 ++++++++++++++---- tests/out/glsl/math-functions.main.Vertex.glsl | 3 +-- tests/out/hlsl/math-functions.hlsl | 3 +-- 4 files changed, 20 insertions(+), 9 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index eed219057b..2004ed5bee 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -1146,7 +1146,10 @@ impl<'a, W: Write> Writer<'a, W> { arg, .. } => { - self.need_bake_expressions.insert(arg); + let inner = info[fun_handle].ty.inner_with(&self.module.types); + if let Some(crate::ScalarKind::Sint) = inner.scalar_kind() { + self.need_bake_expressions.insert(arg); + } } _ => (), } diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index fd9b535796..ff30ba93a4 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -101,7 +101,12 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { /// /// # Notes /// Clears `need_bake_expressions` set before adding to it - fn update_expressions_to_bake(&mut self, func: &crate::Function, info: &valid::FunctionInfo) { + fn update_expressions_to_bake( + &mut self, + module: &Module, + func: &crate::Function, + info: &valid::FunctionInfo, + ) { use crate::Expression; self.need_bake_expressions.clear(); for (fun_handle, expr) in func.expressions.iter() { @@ -116,10 +121,15 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { crate::MathFunction::Asinh | crate::MathFunction::Acosh | crate::MathFunction::Atanh - | crate::MathFunction::Unpack2x16float - | crate::MathFunction::CountLeadingZeros => { + | crate::MathFunction::Unpack2x16float => { self.need_bake_expressions.insert(arg); } + crate::MathFunction::CountLeadingZeros => { + let inner = info[fun_handle].ty.inner_with(&module.types); + if let Some(crate::ScalarKind::Sint) = inner.scalar_kind() { + self.need_bake_expressions.insert(arg); + } + } _ => (), }, _ => (), @@ -1070,7 +1080,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { ) -> BackendResult { // Function Declaration Syntax - https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-function-syntax - self.update_expressions_to_bake(func, info); + self.update_expressions_to_bake(module, func, info); // Write modifier if let Some(crate::FunctionResult { diff --git a/tests/out/glsl/math-functions.main.Vertex.glsl b/tests/out/glsl/math-functions.main.Vertex.glsl index b46e68de88..8cfa6a10b5 100644 --- a/tests/out/glsl/math-functions.main.Vertex.glsl +++ b/tests/out/glsl/math-functions.main.Vertex.glsl @@ -18,7 +18,6 @@ void main() { uint clz_b = uint(31 - findMSB(1u)); ivec2 _e20 = ivec2(-1); ivec2 clz_c = mix(ivec2(31) - findMSB(_e20), ivec2(0), lessThan(_e20, ivec2(0))); - uvec2 _e23 = uvec2(1u); - uvec2 clz_d = uvec2(ivec2(31) - findMSB(_e23)); + uvec2 clz_d = uvec2(ivec2(31) - findMSB(uvec2(1u))); } diff --git a/tests/out/hlsl/math-functions.hlsl b/tests/out/hlsl/math-functions.hlsl index 0b97105598..2a95c849c9 100644 --- a/tests/out/hlsl/math-functions.hlsl +++ b/tests/out/hlsl/math-functions.hlsl @@ -14,6 +14,5 @@ void main() uint clz_b = asuint(31 - firstbithigh(1u)); int2 _expr20 = (-1).xx; int2 clz_c = (_expr20 < (0).xx ? (0).xx : (31).xx - firstbithigh(_expr20)); - uint2 _expr23 = (1u).xx; - uint2 clz_d = asuint((31).xx - firstbithigh(_expr23)); + uint2 clz_d = asuint((31).xx - firstbithigh((1u).xx)); } From 1e08de33200d231e766938a87cffe9eb6366a01b Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Mon, 30 Jan 2023 21:21:15 -0500 Subject: [PATCH 06/10] [glsl-out] CountLeadingZeros without findMSB --- src/back/glsl/mod.rs | 77 +++++++++++-------- .../out/glsl/math-functions.main.Vertex.glsl | 9 ++- 2 files changed, 50 insertions(+), 36 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 2004ed5bee..12d2536258 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -1122,35 +1122,26 @@ impl<'a, W: Write> Writer<'a, W> { } match *expr { - Expression::Math { - fun: crate::MathFunction::Dot, - arg, - arg1, - .. - } => { - // if the expression is a Dot product with integer arguments, - // then the args needs baking as well - let inner = info[fun_handle].ty.inner_with(&self.module.types); - if let TypeInner::Scalar { kind, .. } = *inner { - match kind { - crate::ScalarKind::Sint | crate::ScalarKind::Uint => { - self.need_bake_expressions.insert(arg); - self.need_bake_expressions.insert(arg1.unwrap()); + Expression::Math { fun, arg, arg1, .. } => match fun { + crate::MathFunction::Dot => { + // if the expression is a Dot product with integer arguments, + // then the args needs baking as well + let inner = info[fun_handle].ty.inner_with(&self.module.types); + if let TypeInner::Scalar { kind, .. } = *inner { + match kind { + crate::ScalarKind::Sint | crate::ScalarKind::Uint => { + self.need_bake_expressions.insert(arg); + self.need_bake_expressions.insert(arg1.unwrap()); + } + _ => (), } - _ => (), } } - } - Expression::Math { - fun: crate::MathFunction::CountLeadingZeros, - arg, - .. - } => { - let inner = info[fun_handle].ty.inner_with(&self.module.types); - if let Some(crate::ScalarKind::Sint) = inner.scalar_kind() { + crate::MathFunction::CountLeadingZeros => { self.need_bake_expressions.insert(arg); } - } + _ => (), + }, _ => (), } } @@ -2939,33 +2930,55 @@ impl<'a, W: Write> Writer<'a, W> { Mf::Determinant => "determinant", // bits Mf::CountLeadingZeros => { + use std::f32::consts::LOG2_E; + match *ctx.info[arg].ty.inner_with(&self.module.types) { crate::TypeInner::Vector { size, kind, .. } => { let s = back::vector_size_str(size); + write!(self.out, "mix(")?; + if let crate::ScalarKind::Uint = kind { - write!(self.out, "uvec{s}(ivec{s}(31) - findMSB(")?; + write!(self.out, "uvec{s}(")?; + } else { + write!(self.out, "ivec{s}(")?; + } + + write!(self.out, "vec{s}(31.0) - floor(log(vec{s}(")?; + self.write_expr(arg, ctx)?; + write!(self.out, ") + 0.5) * {LOG2_E})), ")?; + + if let crate::ScalarKind::Uint = kind { + write!(self.out, "mix(uvec{s}(0u), uvec{s}(32u), equal(")?; self.write_expr(arg, ctx)?; - write!(self.out, "))")?; + write!(self.out, ", uvec{s}(0u))), lessThanEqual(")?; + self.write_expr(arg, ctx)?; + write!(self.out, ", uvec{s}(0u)))")?; } else { - write!(self.out, "mix(ivec{s}(31) - findMSB(")?; + write!(self.out, "mix(ivec{s}(0u), ivec{s}(32u), equal(")?; self.write_expr(arg, ctx)?; - write!(self.out, "), ivec{s}(0), lessThan(")?; + write!(self.out, ", ivec{s}(0))), lessThanEqual(")?; self.write_expr(arg, ctx)?; write!(self.out, ", ivec{s}(0)))")?; } } crate::TypeInner::Scalar { kind, .. } => { + write!(self.out, "(")?; + self.write_expr(arg, ctx)?; + if let crate::ScalarKind::Uint = kind { - write!(self.out, "uint(31 - findMSB(")?; + write!(self.out, " <= 0u ? (")?; + self.write_expr(arg, ctx)?; + write!(self.out, " == 0u ? 32u : 0u) : uint(")?; } else { - write!(self.out, "(")?; + write!(self.out, " <= 0 ? (")?; self.write_expr(arg, ctx)?; - write!(self.out, " < 0 ? 0 : 31 - findMSB(")?; + write!(self.out, " == 0 ? 32 : 0) : int(")?; } + write!(self.out, "31.0 - floor(log(float(")?; self.write_expr(arg, ctx)?; - write!(self.out, "))")?; + write!(self.out, ") + 0.5) * {LOG2_E})))")?; } _ => unreachable!(), }; diff --git a/tests/out/glsl/math-functions.main.Vertex.glsl b/tests/out/glsl/math-functions.main.Vertex.glsl index 8cfa6a10b5..0cb4130f3e 100644 --- a/tests/out/glsl/math-functions.main.Vertex.glsl +++ b/tests/out/glsl/math-functions.main.Vertex.glsl @@ -14,10 +14,11 @@ void main() { vec4 g = refract(v, v, 1.0); int const_dot = ( + ivec2(0, 0).x * ivec2(0, 0).x + ivec2(0, 0).y * ivec2(0, 0).y); uint first_leading_bit_abs = uint(findMSB(uint(abs(int(0u))))); - int clz_a = (-1 < 0 ? 0 : 31 - findMSB(-1)); - uint clz_b = uint(31 - findMSB(1u)); + int clz_a = (-1 <= 0 ? (-1 == 0 ? 32 : 0) : int(31.0 - floor(log(float(-1) + 0.5) * 1.442695))); + uint clz_b = (1u <= 0u ? (1u == 0u ? 32u : 0u) : uint(31.0 - floor(log(float(1u) + 0.5) * 1.442695))); ivec2 _e20 = ivec2(-1); - ivec2 clz_c = mix(ivec2(31) - findMSB(_e20), ivec2(0), lessThan(_e20, ivec2(0))); - uvec2 clz_d = uvec2(ivec2(31) - findMSB(uvec2(1u))); + ivec2 clz_c = mix(ivec2(vec2(31.0) - floor(log(vec2(_e20) + 0.5) * 1.442695)), mix(ivec2(0u), ivec2(32u), equal(_e20, ivec2(0))), lessThanEqual(_e20, ivec2(0))); + uvec2 _e23 = uvec2(1u); + uvec2 clz_d = mix(uvec2(vec2(31.0) - floor(log(vec2(_e23) + 0.5) * 1.442695)), mix(uvec2(0u), uvec2(32u), equal(_e23, uvec2(0u))), lessThanEqual(_e23, uvec2(0u))); } From c49fe99d9d02ae98cc443fc5fb753ff9b57e34e7 Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Tue, 31 Jan 2023 01:27:04 -0500 Subject: [PATCH 07/10] Don't check negatives when uint --- src/back/glsl/mod.rs | 8 ++------ tests/out/glsl/math-functions.main.Vertex.glsl | 4 ++-- 2 files changed, 4 insertions(+), 8 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 12d2536258..41f93b90ba 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -2949,9 +2949,7 @@ impl<'a, W: Write> Writer<'a, W> { write!(self.out, ") + 0.5) * {LOG2_E})), ")?; if let crate::ScalarKind::Uint = kind { - write!(self.out, "mix(uvec{s}(0u), uvec{s}(32u), equal(")?; - self.write_expr(arg, ctx)?; - write!(self.out, ", uvec{s}(0u))), lessThanEqual(")?; + write!(self.out, "uvec{s}(32u), lessThanEqual(")?; self.write_expr(arg, ctx)?; write!(self.out, ", uvec{s}(0u)))")?; } else { @@ -2967,9 +2965,7 @@ impl<'a, W: Write> Writer<'a, W> { self.write_expr(arg, ctx)?; if let crate::ScalarKind::Uint = kind { - write!(self.out, " <= 0u ? (")?; - self.write_expr(arg, ctx)?; - write!(self.out, " == 0u ? 32u : 0u) : uint(")?; + write!(self.out, " == 0u ? 32u : uint(")?; } else { write!(self.out, " <= 0 ? (")?; self.write_expr(arg, ctx)?; diff --git a/tests/out/glsl/math-functions.main.Vertex.glsl b/tests/out/glsl/math-functions.main.Vertex.glsl index 0cb4130f3e..f08e6292c2 100644 --- a/tests/out/glsl/math-functions.main.Vertex.glsl +++ b/tests/out/glsl/math-functions.main.Vertex.glsl @@ -15,10 +15,10 @@ void main() { int const_dot = ( + ivec2(0, 0).x * ivec2(0, 0).x + ivec2(0, 0).y * ivec2(0, 0).y); uint first_leading_bit_abs = uint(findMSB(uint(abs(int(0u))))); int clz_a = (-1 <= 0 ? (-1 == 0 ? 32 : 0) : int(31.0 - floor(log(float(-1) + 0.5) * 1.442695))); - uint clz_b = (1u <= 0u ? (1u == 0u ? 32u : 0u) : uint(31.0 - floor(log(float(1u) + 0.5) * 1.442695))); + uint clz_b = (1u == 0u ? 32u : uint(31.0 - floor(log(float(1u) + 0.5) * 1.442695))); ivec2 _e20 = ivec2(-1); ivec2 clz_c = mix(ivec2(vec2(31.0) - floor(log(vec2(_e20) + 0.5) * 1.442695)), mix(ivec2(0u), ivec2(32u), equal(_e20, ivec2(0))), lessThanEqual(_e20, ivec2(0))); uvec2 _e23 = uvec2(1u); - uvec2 clz_d = mix(uvec2(vec2(31.0) - floor(log(vec2(_e23) + 0.5) * 1.442695)), mix(uvec2(0u), uvec2(32u), equal(_e23, uvec2(0u))), lessThanEqual(_e23, uvec2(0u))); + uvec2 clz_d = mix(uvec2(vec2(31.0) - floor(log(vec2(_e23) + 0.5) * 1.442695)), uvec2(32u), lessThanEqual(_e23, uvec2(0u))); } From 05f264968fc7027ae4b54d4d18182a8ca1d02a8a Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Tue, 31 Jan 2023 02:48:15 -0500 Subject: [PATCH 08/10] Perform the type conv after mix --- src/back/glsl/mod.rs | 14 ++++++-------- tests/out/glsl/math-functions.main.Vertex.glsl | 4 ++-- 2 files changed, 8 insertions(+), 10 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 41f93b90ba..89a4a69f74 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -2936,28 +2936,26 @@ impl<'a, W: Write> Writer<'a, W> { crate::TypeInner::Vector { size, kind, .. } => { let s = back::vector_size_str(size); - write!(self.out, "mix(")?; - if let crate::ScalarKind::Uint = kind { write!(self.out, "uvec{s}(")?; } else { write!(self.out, "ivec{s}(")?; } - write!(self.out, "vec{s}(31.0) - floor(log(vec{s}(")?; + write!(self.out, "mix(vec{s}(31.0) - floor(log(vec{s}(")?; self.write_expr(arg, ctx)?; - write!(self.out, ") + 0.5) * {LOG2_E})), ")?; + write!(self.out, ") + 0.5) * {LOG2_E}), ")?; if let crate::ScalarKind::Uint = kind { - write!(self.out, "uvec{s}(32u), lessThanEqual(")?; + write!(self.out, "vec{s}(32.0), lessThanEqual(")?; self.write_expr(arg, ctx)?; - write!(self.out, ", uvec{s}(0u)))")?; + write!(self.out, ", uvec{s}(0u))))")?; } else { - write!(self.out, "mix(ivec{s}(0u), ivec{s}(32u), equal(")?; + write!(self.out, "mix(vec{s}(0.0), vec{s}(32.0), equal(")?; self.write_expr(arg, ctx)?; write!(self.out, ", ivec{s}(0))), lessThanEqual(")?; self.write_expr(arg, ctx)?; - write!(self.out, ", ivec{s}(0)))")?; + write!(self.out, ", ivec{s}(0))))")?; } } crate::TypeInner::Scalar { kind, .. } => { diff --git a/tests/out/glsl/math-functions.main.Vertex.glsl b/tests/out/glsl/math-functions.main.Vertex.glsl index f08e6292c2..a401138ec7 100644 --- a/tests/out/glsl/math-functions.main.Vertex.glsl +++ b/tests/out/glsl/math-functions.main.Vertex.glsl @@ -17,8 +17,8 @@ void main() { int clz_a = (-1 <= 0 ? (-1 == 0 ? 32 : 0) : int(31.0 - floor(log(float(-1) + 0.5) * 1.442695))); uint clz_b = (1u == 0u ? 32u : uint(31.0 - floor(log(float(1u) + 0.5) * 1.442695))); ivec2 _e20 = ivec2(-1); - ivec2 clz_c = mix(ivec2(vec2(31.0) - floor(log(vec2(_e20) + 0.5) * 1.442695)), mix(ivec2(0u), ivec2(32u), equal(_e20, ivec2(0))), lessThanEqual(_e20, ivec2(0))); + ivec2 clz_c = ivec2(mix(vec2(31.0) - floor(log(vec2(_e20) + 0.5) * 1.442695), mix(vec2(0.0), vec2(32.0), equal(_e20, ivec2(0))), lessThanEqual(_e20, ivec2(0)))); uvec2 _e23 = uvec2(1u); - uvec2 clz_d = mix(uvec2(vec2(31.0) - floor(log(vec2(_e23) + 0.5) * 1.442695)), uvec2(32u), lessThanEqual(_e23, uvec2(0u))); + uvec2 clz_d = uvec2(mix(vec2(31.0) - floor(log(vec2(_e23) + 0.5) * 1.442695), vec2(32.0), lessThanEqual(_e23, uvec2(0u)))); } From 254e806edbdac6b37b06fb373c20bd02ec441474 Mon Sep 17 00:00:00 2001 From: teoxoy <28601907+teoxoy@users.noreply.github.com> Date: Tue, 31 Jan 2023 21:07:01 +0100 Subject: [PATCH 09/10] use log2 --- src/back/glsl/mod.rs | 10 ++++------ tests/out/glsl/math-functions.main.Vertex.glsl | 8 ++++---- 2 files changed, 8 insertions(+), 10 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 89a4a69f74..d975c5d65a 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -2930,8 +2930,6 @@ impl<'a, W: Write> Writer<'a, W> { Mf::Determinant => "determinant", // bits Mf::CountLeadingZeros => { - use std::f32::consts::LOG2_E; - match *ctx.info[arg].ty.inner_with(&self.module.types) { crate::TypeInner::Vector { size, kind, .. } => { let s = back::vector_size_str(size); @@ -2942,9 +2940,9 @@ impl<'a, W: Write> Writer<'a, W> { write!(self.out, "ivec{s}(")?; } - write!(self.out, "mix(vec{s}(31.0) - floor(log(vec{s}(")?; + write!(self.out, "mix(vec{s}(31.0) - floor(log2(vec{s}(")?; self.write_expr(arg, ctx)?; - write!(self.out, ") + 0.5) * {LOG2_E}), ")?; + write!(self.out, ") + 0.5)), ")?; if let crate::ScalarKind::Uint = kind { write!(self.out, "vec{s}(32.0), lessThanEqual(")?; @@ -2970,9 +2968,9 @@ impl<'a, W: Write> Writer<'a, W> { write!(self.out, " == 0 ? 32 : 0) : int(")?; } - write!(self.out, "31.0 - floor(log(float(")?; + write!(self.out, "31.0 - floor(log2(float(")?; self.write_expr(arg, ctx)?; - write!(self.out, ") + 0.5) * {LOG2_E})))")?; + write!(self.out, ") + 0.5))))")?; } _ => unreachable!(), }; diff --git a/tests/out/glsl/math-functions.main.Vertex.glsl b/tests/out/glsl/math-functions.main.Vertex.glsl index a401138ec7..d863ada72b 100644 --- a/tests/out/glsl/math-functions.main.Vertex.glsl +++ b/tests/out/glsl/math-functions.main.Vertex.glsl @@ -14,11 +14,11 @@ void main() { vec4 g = refract(v, v, 1.0); int const_dot = ( + ivec2(0, 0).x * ivec2(0, 0).x + ivec2(0, 0).y * ivec2(0, 0).y); uint first_leading_bit_abs = uint(findMSB(uint(abs(int(0u))))); - int clz_a = (-1 <= 0 ? (-1 == 0 ? 32 : 0) : int(31.0 - floor(log(float(-1) + 0.5) * 1.442695))); - uint clz_b = (1u == 0u ? 32u : uint(31.0 - floor(log(float(1u) + 0.5) * 1.442695))); + int clz_a = (-1 <= 0 ? (-1 == 0 ? 32 : 0) : int(31.0 - floor(log2(float(-1) + 0.5)))); + uint clz_b = (1u == 0u ? 32u : uint(31.0 - floor(log2(float(1u) + 0.5)))); ivec2 _e20 = ivec2(-1); - ivec2 clz_c = ivec2(mix(vec2(31.0) - floor(log(vec2(_e20) + 0.5) * 1.442695), mix(vec2(0.0), vec2(32.0), equal(_e20, ivec2(0))), lessThanEqual(_e20, ivec2(0)))); + ivec2 clz_c = ivec2(mix(vec2(31.0) - floor(log2(vec2(_e20) + 0.5)), mix(vec2(0.0), vec2(32.0), equal(_e20, ivec2(0))), lessThanEqual(_e20, ivec2(0)))); uvec2 _e23 = uvec2(1u); - uvec2 clz_d = uvec2(mix(vec2(31.0) - floor(log(vec2(_e23) + 0.5) * 1.442695), vec2(32.0), lessThanEqual(_e23, uvec2(0u)))); + uvec2 clz_d = uvec2(mix(vec2(31.0) - floor(log2(vec2(_e23) + 0.5)), vec2(32.0), lessThanEqual(_e23, uvec2(0u)))); } From c66ddba667bfdf0d1a9a0841cf6f9a91b42fced7 Mon Sep 17 00:00:00 2001 From: teoxoy <28601907+teoxoy@users.noreply.github.com> Date: Tue, 31 Jan 2023 21:21:03 +0100 Subject: [PATCH 10/10] fix clippy lints --- src/back/glsl/mod.rs | 11 +++++------ src/back/hlsl/writer.rs | 9 ++++----- 2 files changed, 9 insertions(+), 11 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index d975c5d65a..6c90fcb26d 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -1121,8 +1121,8 @@ impl<'a, W: Write> Writer<'a, W> { self.need_bake_expressions.insert(fun_handle); } - match *expr { - Expression::Math { fun, arg, arg1, .. } => match fun { + if let Expression::Math { fun, arg, arg1, .. } = *expr { + match fun { crate::MathFunction::Dot => { // if the expression is a Dot product with integer arguments, // then the args needs baking as well @@ -1133,16 +1133,15 @@ impl<'a, W: Write> Writer<'a, W> { self.need_bake_expressions.insert(arg); self.need_bake_expressions.insert(arg1.unwrap()); } - _ => (), + _ => {} } } } crate::MathFunction::CountLeadingZeros => { self.need_bake_expressions.insert(arg); } - _ => (), - }, - _ => (), + _ => {} + } } } } diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index ff30ba93a4..2af7a3524b 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -116,8 +116,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { self.need_bake_expressions.insert(fun_handle); } - match *expr { - Expression::Math { fun, arg, .. } => match fun { + if let Expression::Math { fun, arg, .. } = *expr { + match fun { crate::MathFunction::Asinh | crate::MathFunction::Acosh | crate::MathFunction::Atanh @@ -130,9 +130,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { self.need_bake_expressions.insert(arg); } } - _ => (), - }, - _ => (), + _ => {} + } } } }