From 336b890eb5eb13e9034f932fc0a00a47a40c7f4c Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Wed, 10 Aug 2022 17:06:31 -0700 Subject: [PATCH 1/6] adds saturate to supporting targets --- src/back/glsl/mod.rs | 1 + src/back/hlsl/writer.rs | 1 + src/back/msl/writer.rs | 1 + src/back/spv/block.rs | 1 + src/back/wgsl/writer.rs | 1 + src/front/wgsl/conv.rs | 1 + src/lib.rs | 1 + src/proc/mod.rs | 1 + src/proc/typifier.rs | 1 + src/valid/expression.rs | 14 ++++++++++++++ 10 files changed, 23 insertions(+) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index e714b64475..d1a5a589a7 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -2749,6 +2749,7 @@ impl<'a, W: Write> Writer<'a, W> { Mf::Min => "min", Mf::Max => "max", Mf::Clamp => "clamp", + Mf::Saturate => todo!(), // trigonometry Mf::Cos => "cos", Mf::Cosh => "cosh", diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index a37a89444a..ae761ead4f 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -2433,6 +2433,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { Mf::Min => Function::Regular("min"), Mf::Max => Function::Regular("max"), Mf::Clamp => Function::Regular("clamp"), + Mf::Saturate => Function::Regular("saturate"), // trigonometry Mf::Cos => Function::Regular("cos"), Mf::Cosh => Function::Regular("cosh"), diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index c35786f7ad..9a59289bef 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -1628,6 +1628,7 @@ impl Writer { Mf::Min => "min", Mf::Max => "max", Mf::Clamp => "clamp", + Mf::Saturate => "saturate", // trigonometry Mf::Cos => "cos", Mf::Cosh => "cosh", diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index 1c475e49b0..3ad51979c0 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -696,6 +696,7 @@ impl<'w> BlockContext<'w> { Some(crate::ScalarKind::Uint) => spirv::GLOp::UClamp, other => unimplemented!("Unexpected max({:?})", other), }), + Mf::Saturate => todo!(), // trigonometry Mf::Sin => MathOp::Ext(spirv::GLOp::Sin), Mf::Sinh => MathOp::Ext(spirv::GLOp::Sinh), diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index 8a7a71cdc2..0be7ef072a 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -1511,6 +1511,7 @@ impl Writer { Mf::Min => Function::Regular("min"), Mf::Max => Function::Regular("max"), Mf::Clamp => Function::Regular("clamp"), + Mf::Saturate => Function::Regular("saturate"), // trigonometry Mf::Cos => Function::Regular("cos"), Mf::Cosh => Function::Regular("cosh"), diff --git a/src/front/wgsl/conv.rs b/src/front/wgsl/conv.rs index c1bbb89715..ba41648757 100644 --- a/src/front/wgsl/conv.rs +++ b/src/front/wgsl/conv.rs @@ -140,6 +140,7 @@ pub fn map_standard_fun(word: &str) -> Option { "min" => Mf::Min, "max" => Mf::Max, "clamp" => Mf::Clamp, + "saturate" => Mf::Saturate, // trigonometry "cos" => Mf::Cos, "cosh" => Mf::Cosh, diff --git a/src/lib.rs b/src/lib.rs index 9f45c082aa..107534c8b0 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -952,6 +952,7 @@ pub enum MathFunction { Min, Max, Clamp, + Saturate, // trigonometry Cos, Cosh, diff --git a/src/proc/mod.rs b/src/proc/mod.rs index 823a190dec..0b9f406af4 100644 --- a/src/proc/mod.rs +++ b/src/proc/mod.rs @@ -228,6 +228,7 @@ impl super::MathFunction { Self::Min => 2, Self::Max => 2, Self::Clamp => 3, + Self::Saturate => 1, // trigonometry Self::Cos => 1, Self::Cosh => 1, diff --git a/src/proc/typifier.rs b/src/proc/typifier.rs index a5304b738c..9a5922ea76 100644 --- a/src/proc/typifier.rs +++ b/src/proc/typifier.rs @@ -704,6 +704,7 @@ impl<'a> ResolveContext<'a> { Mf::Min | Mf::Max | Mf::Clamp | + Mf::Saturate | // trigonometry Mf::Cos | Mf::Cosh | diff --git a/src/valid/expression.rs b/src/valid/expression.rs index f433b456b1..18ee9dd9b1 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -1040,6 +1040,20 @@ impl super::Validator { )); } } + Mf::Saturate => { + if arg1_ty.is_some() | arg2_ty.is_some() | arg3_ty.is_some() { + return Err(ExpressionError::WrongArgumentCount(fun)); + } + match *arg_ty { + Ti::Scalar { + kind: Sk::Float, .. + } + | Ti::Vector { + kind: Sk::Float, .. + } => {} + _ => return Err(ExpressionError::InvalidArgumentType(fun, 0, arg)), + } + } Mf::Cos | Mf::Cosh | Mf::Sin From 12ff1c7c2a67c67038cb496e0180b9c06ae2525b Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Wed, 10 Aug 2022 17:55:45 -0700 Subject: [PATCH 2/6] uses clamp in place of saturate when spv --- src/back/spv/block.rs | 83 ++++++++++++++++++++++++++++++++++++++++++- 1 file changed, 82 insertions(+), 1 deletion(-) diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index 3ad51979c0..5490037127 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -696,7 +696,88 @@ impl<'w> BlockContext<'w> { Some(crate::ScalarKind::Uint) => spirv::GLOp::UClamp, other => unimplemented!("Unexpected max({:?})", other), }), - Mf::Saturate => todo!(), + Mf::Saturate => match *arg_ty { + crate::TypeInner::Vector { + size, + kind: crate::ScalarKind::Float, + width, + } => { + let result_type_id = { + let value = LocalType::Value { + vector_size: Some(size), + kind: crate::ScalarKind::Float, + width, + pointer_space: None, + }; + + self.get_type_id(LookupType::Local(value)) + }; + + let arg1_id = { + self.temp_list.clear(); + self.temp_list.resize( + size as _, + self.writer + .get_constant_scalar(crate::ScalarValue::Float(0.0), width), + ); + + let id = self.gen_id(); + block.body.push(Instruction::composite_construct( + result_type_id, + id, + &self.temp_list, + )); + + id + }; + + let arg2_id = { + self.temp_list.clear(); + self.temp_list.resize( + size as _, + self.writer + .get_constant_scalar(crate::ScalarValue::Float(1.0), width), + ); + + let id = self.gen_id(); + block.body.push(Instruction::composite_construct( + result_type_id, + id, + &self.temp_list, + )); + + id + }; + + MathOp::Custom(Instruction::ext_inst( + self.writer.gl450_ext_inst_id, + spirv::GLOp::FClamp, + result_type_id, + id, + &[arg0_id, arg1_id, arg2_id], + )) + } + crate::TypeInner::Scalar { + kind: crate::ScalarKind::Float, + width, + } => { + let arg1_id = self + .writer + .get_constant_scalar(crate::ScalarValue::Float(0.0), width); + let arg2_id = self + .writer + .get_constant_scalar(crate::ScalarValue::Float(1.0), width); + + MathOp::Custom(Instruction::ext_inst( + self.writer.gl450_ext_inst_id, + spirv::GLOp::FClamp, + result_type_id, + id, + &[arg0_id, arg1_id, arg2_id], + )) + } + ref other => unimplemented!("Unexpected saturate({:?})", other), + }, // trigonometry Mf::Sin => MathOp::Ext(spirv::GLOp::Sin), Mf::Sinh => MathOp::Ext(spirv::GLOp::Sinh), From 6766aa7ef2f9a3619c49a6ba3adaa1e96c09192f Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Wed, 10 Aug 2022 18:04:13 -0700 Subject: [PATCH 3/6] uses clamp in place of saturate when glsl --- src/back/glsl/mod.rs | 27 ++++++++++++++++++++++++++- 1 file changed, 26 insertions(+), 1 deletion(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index d1a5a589a7..c8a9d7059d 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -2749,7 +2749,32 @@ impl<'a, W: Write> Writer<'a, W> { Mf::Min => "min", Mf::Max => "max", Mf::Clamp => "clamp", - Mf::Saturate => todo!(), + Mf::Saturate => { + write!(self.out, "clamp(")?; + + self.write_expr(arg, ctx)?; + + match ctx.info[arg].ty.inner_with(&self.module.types) { + crate::TypeInner::Vector { + size, + kind: crate::ScalarKind::Float, + .. + } => write!( + self.out, + ", vec{}(0.0), vec{0}(1.0)", + back::vector_size_str(*size) + )?, + crate::TypeInner::Scalar { + kind: crate::ScalarKind::Float, + .. + } => write!(self.out, ", 0.0, 1.0")?, + other => unreachable!("Unexpected type {:?}", other), + } + + write!(self.out, ")")?; + + return Ok(()); + } // trigonometry Mf::Cos => "cos", Mf::Cosh => "cosh", From 4c9772c5b0261430817e00e41a569eb8db83e51c Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Wed, 10 Aug 2022 18:35:07 -0700 Subject: [PATCH 4/6] derefs the match expr --- src/back/glsl/mod.rs | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index c8a9d7059d..3f4566dbfa 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -2754,7 +2754,7 @@ impl<'a, W: Write> Writer<'a, W> { self.write_expr(arg, ctx)?; - match ctx.info[arg].ty.inner_with(&self.module.types) { + match *ctx.info[arg].ty.inner_with(&self.module.types) { crate::TypeInner::Vector { size, kind: crate::ScalarKind::Float, @@ -2762,13 +2762,13 @@ impl<'a, W: Write> Writer<'a, W> { } => write!( self.out, ", vec{}(0.0), vec{0}(1.0)", - back::vector_size_str(*size) + back::vector_size_str(size) )?, crate::TypeInner::Scalar { kind: crate::ScalarKind::Float, .. } => write!(self.out, ", 0.0, 1.0")?, - other => unreachable!("Unexpected type {:?}", other), + ref other => unreachable!("Unexpected type {:?}", other), } write!(self.out, ")")?; From 76807ba99eae8d4f7f122b5fb3cb46701a939563 Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins <85699459+evahop@users.noreply.github.com> Date: Thu, 11 Aug 2022 14:49:59 -0700 Subject: [PATCH 5/6] simplifies spv block MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-authored-by: João Capucho --- src/back/spv/block.rs | 131 +++++++++++++++++------------------------- 1 file changed, 54 insertions(+), 77 deletions(-) diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index 5490037127..af2326ac14 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -696,88 +696,65 @@ impl<'w> BlockContext<'w> { Some(crate::ScalarKind::Uint) => spirv::GLOp::UClamp, other => unimplemented!("Unexpected max({:?})", other), }), - Mf::Saturate => match *arg_ty { - crate::TypeInner::Vector { - size, + Mf::Saturate => { + let (maybe_size, width) = match *arg_ty { + crate::TypeInner::Vector { + size, + width, + .. + } => (Some(size), width), + crate::TypeInner::Scalar { width, .. } => (None, width), + ref other => unimplemented!("Unexpected saturate({:?})", other), + }; + + let mut arg1_id = self + .writer + .get_constant_scalar(crate::ScalarValue::Float(0.0), width); + let mut arg2_id = self + .writer + .get_constant_scalar(crate::ScalarValue::Float(1.0), width); + + if let Some(size) = maybe_size { + let value = LocalType::Value { + vector_size: Some(size), kind: crate::ScalarKind::Float, width, - } => { - let result_type_id = { - let value = LocalType::Value { - vector_size: Some(size), - kind: crate::ScalarKind::Float, - width, - pointer_space: None, - }; - - self.get_type_id(LookupType::Local(value)) - }; - - let arg1_id = { - self.temp_list.clear(); - self.temp_list.resize( - size as _, - self.writer - .get_constant_scalar(crate::ScalarValue::Float(0.0), width), - ); - - let id = self.gen_id(); - block.body.push(Instruction::composite_construct( - result_type_id, - id, - &self.temp_list, - )); - - id - }; - - let arg2_id = { - self.temp_list.clear(); - self.temp_list.resize( - size as _, - self.writer - .get_constant_scalar(crate::ScalarValue::Float(1.0), width), - ); + pointer_space: None, + }; - let id = self.gen_id(); - block.body.push(Instruction::composite_construct( - result_type_id, - id, - &self.temp_list, - )); + let result_type_id = self.get_type_id(LookupType::Local(value)); + + self.temp_list.clear(); + self.temp_list.resize(size as _, arg1_id); - id - }; + let id = self.gen_id(); + block.body.push(Instruction::composite_construct( + result_type_id, + id, + &self.temp_list, + )); + arg1_id = id; + + self.temp_list.clear(); + self.temp_list.resize(size as _, arg2_id); - MathOp::Custom(Instruction::ext_inst( - self.writer.gl450_ext_inst_id, - spirv::GLOp::FClamp, - result_type_id, - id, - &[arg0_id, arg1_id, arg2_id], - )) - } - crate::TypeInner::Scalar { - kind: crate::ScalarKind::Float, - width, - } => { - let arg1_id = self - .writer - .get_constant_scalar(crate::ScalarValue::Float(0.0), width); - let arg2_id = self - .writer - .get_constant_scalar(crate::ScalarValue::Float(1.0), width); - - MathOp::Custom(Instruction::ext_inst( - self.writer.gl450_ext_inst_id, - spirv::GLOp::FClamp, - result_type_id, - id, - &[arg0_id, arg1_id, arg2_id], - )) - } - ref other => unimplemented!("Unexpected saturate({:?})", other), - }, + let id = self.gen_id(); + block.body.push(Instruction::composite_construct( + result_type_id, + id, + &self.temp_list, + )); + arg2_id = id; + } + + MathOp::Custom(Instruction::ext_inst( + self.writer.gl450_ext_inst_id, + spirv::GLOp::FClamp, + result_type_id, + id, + &[arg0_id, arg1_id, arg2_id], + )) + }, // trigonometry Mf::Sin => MathOp::Ext(spirv::GLOp::Sin), Mf::Sinh => MathOp::Ext(spirv::GLOp::Sinh), From 972bc8a82a2349dd4f7bbdfbc6027e0820ad2a08 Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Thu, 11 Aug 2022 15:22:53 -0700 Subject: [PATCH 6/6] makes suggested changes --- src/back/glsl/mod.rs | 12 +-- src/back/spv/block.rs | 100 +++++++++--------- src/valid/expression.rs | 17 +-- tests/in/math-functions.wgsl | 1 + .../out/glsl/math-functions.main.Vertex.glsl | 1 + tests/out/hlsl/math-functions.hlsl | 1 + tests/out/msl/math-functions.msl | 1 + tests/out/spv/math-functions.spvasm | 27 ++--- tests/out/wgsl/math-functions.wgsl | 1 + 9 files changed, 72 insertions(+), 89 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 3f4566dbfa..a3f2a53836 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -2755,20 +2755,12 @@ impl<'a, W: Write> Writer<'a, W> { self.write_expr(arg, ctx)?; match *ctx.info[arg].ty.inner_with(&self.module.types) { - crate::TypeInner::Vector { - size, - kind: crate::ScalarKind::Float, - .. - } => write!( + crate::TypeInner::Vector { size, .. } => write!( self.out, ", vec{}(0.0), vec{0}(1.0)", back::vector_size_str(size) )?, - crate::TypeInner::Scalar { - kind: crate::ScalarKind::Float, - .. - } => write!(self.out, ", 0.0, 1.0")?, - ref other => unreachable!("Unexpected type {:?}", other), + _ => write!(self.out, ", 0.0, 1.0")?, } write!(self.out, ")")?; diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index af2326ac14..8a388be22e 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -697,64 +697,60 @@ impl<'w> BlockContext<'w> { other => unimplemented!("Unexpected max({:?})", other), }), Mf::Saturate => { - let (maybe_size, width) = match *arg_ty { - crate::TypeInner::Vector { - size, - width, - .. - } => (Some(size), width), - crate::TypeInner::Scalar { width, .. } => (None, width), - ref other => unimplemented!("Unexpected saturate({:?})", other), - }; - - let mut arg1_id = self - .writer - .get_constant_scalar(crate::ScalarValue::Float(0.0), width); - let mut arg2_id = self - .writer - .get_constant_scalar(crate::ScalarValue::Float(1.0), width); - - if let Some(size) = maybe_size { - let value = LocalType::Value { - vector_size: Some(size), - kind: crate::ScalarKind::Float, - width, - pointer_space: None, + let (maybe_size, width) = match *arg_ty { + crate::TypeInner::Vector { size, width, .. } => (Some(size), width), + crate::TypeInner::Scalar { width, .. } => (None, width), + ref other => unimplemented!("Unexpected saturate({:?})", other), }; - let result_type_id = self.get_type_id(LookupType::Local(value)); - - self.temp_list.clear(); - self.temp_list.resize(size as _, arg1_id); + let mut arg1_id = self + .writer + .get_constant_scalar(crate::ScalarValue::Float(0.0), width); + let mut arg2_id = self + .writer + .get_constant_scalar(crate::ScalarValue::Float(1.0), width); + + if let Some(size) = maybe_size { + let value = LocalType::Value { + vector_size: Some(size), + kind: crate::ScalarKind::Float, + width, + pointer_space: None, + }; + + let result_type_id = self.get_type_id(LookupType::Local(value)); - let id = self.gen_id(); - block.body.push(Instruction::composite_construct( - result_type_id, - id, - &self.temp_list, - )); - arg1_id = id; - - self.temp_list.clear(); - self.temp_list.resize(size as _, arg2_id); + self.temp_list.clear(); + self.temp_list.resize(size as _, arg1_id); - let id = self.gen_id(); - block.body.push(Instruction::composite_construct( + let id = self.gen_id(); + block.body.push(Instruction::composite_construct( + result_type_id, + id, + &self.temp_list, + )); + arg1_id = id; + + self.temp_list.clear(); + self.temp_list.resize(size as _, arg2_id); + + let id = self.gen_id(); + block.body.push(Instruction::composite_construct( + result_type_id, + id, + &self.temp_list, + )); + arg2_id = id; + } + + MathOp::Custom(Instruction::ext_inst( + self.writer.gl450_ext_inst_id, + spirv::GLOp::FClamp, result_type_id, id, - &self.temp_list, - )); - arg2_id = id; - } - - MathOp::Custom(Instruction::ext_inst( - self.writer.gl450_ext_inst_id, - spirv::GLOp::FClamp, - result_type_id, - id, - &[arg0_id, arg1_id, arg2_id], - )) - }, + &[arg0_id, arg1_id, arg2_id], + )) + } // trigonometry Mf::Sin => MathOp::Ext(spirv::GLOp::Sin), Mf::Sinh => MathOp::Ext(spirv::GLOp::Sinh), diff --git a/src/valid/expression.rs b/src/valid/expression.rs index 18ee9dd9b1..6279ed3622 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -1040,21 +1040,8 @@ impl super::Validator { )); } } - Mf::Saturate => { - if arg1_ty.is_some() | arg2_ty.is_some() | arg3_ty.is_some() { - return Err(ExpressionError::WrongArgumentCount(fun)); - } - match *arg_ty { - Ti::Scalar { - kind: Sk::Float, .. - } - | Ti::Vector { - kind: Sk::Float, .. - } => {} - _ => return Err(ExpressionError::InvalidArgumentType(fun, 0, arg)), - } - } - Mf::Cos + Mf::Saturate + | Mf::Cos | Mf::Cosh | Mf::Sin | Mf::Sinh diff --git a/tests/in/math-functions.wgsl b/tests/in/math-functions.wgsl index a2e097a04d..9ecac80499 100644 --- a/tests/in/math-functions.wgsl +++ b/tests/in/math-functions.wgsl @@ -6,6 +6,7 @@ fn main() { let b = radians(f); let c = degrees(v); let d = radians(v); + let e = saturate(v); let const_dot = dot(vec2(), vec2()); let first_leading_bit_abs = firstLeadingBit(abs(0u)); } diff --git a/tests/out/glsl/math-functions.main.Vertex.glsl b/tests/out/glsl/math-functions.main.Vertex.glsl index 8d72868299..f0954772e6 100644 --- a/tests/out/glsl/math-functions.main.Vertex.glsl +++ b/tests/out/glsl/math-functions.main.Vertex.glsl @@ -10,6 +10,7 @@ void main() { float b = radians(1.0); vec4 c = degrees(v); vec4 d = radians(v); + vec4 e = clamp(v, vec4(0.0), vec4(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))))); } diff --git a/tests/out/hlsl/math-functions.hlsl b/tests/out/hlsl/math-functions.hlsl index 179cfee66c..827da71b4a 100644 --- a/tests/out/hlsl/math-functions.hlsl +++ b/tests/out/hlsl/math-functions.hlsl @@ -6,6 +6,7 @@ void main() float b = radians(1.0); float4 c = degrees(v); float4 d = radians(v); + float4 e = saturate(v); int const_dot = dot(int2(0, 0), int2(0, 0)); uint first_leading_bit_abs = firstbithigh(abs(0u)); } diff --git a/tests/out/msl/math-functions.msl b/tests/out/msl/math-functions.msl index 87de8418ce..6ee0093d00 100644 --- a/tests/out/msl/math-functions.msl +++ b/tests/out/msl/math-functions.msl @@ -13,6 +13,7 @@ vertex void main_( float b = ((1.0) * 0.017453292519943295474); metal::float4 c = ((v) * 57.295779513082322865); metal::float4 d = ((v) * 0.017453292519943295474); + metal::float4 e = metal::saturate(v); int const_dot = ( + const_type.x * const_type.x + const_type.y * const_type.y); uint first_leading_bit_abs = (((metal::clz(metal::abs(0u)) + 1) % 33) - 1); } diff --git a/tests/out/spv/math-functions.spvasm b/tests/out/spv/math-functions.spvasm index 07a16f754c..fc970f8631 100644 --- a/tests/out/spv/math-functions.spvasm +++ b/tests/out/spv/math-functions.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 33 +; Bound: 36 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 @@ -18,7 +18,7 @@ OpEntryPoint Vertex %13 "main" %11 = OpConstantComposite %10 %6 %6 %14 = OpTypeFunction %2 %16 = OpTypeVector %4 4 -%23 = OpConstantNull %7 +%26 = OpConstantNull %7 %13 = OpFunction %2 None %14 %12 = OpLabel OpBranch %15 @@ -28,15 +28,18 @@ OpBranch %15 %19 = OpExtInst %4 %1 Radians %3 %20 = OpExtInst %16 %1 Degrees %17 %21 = OpExtInst %16 %1 Radians %17 -%24 = OpCompositeExtract %7 %11 0 -%25 = OpCompositeExtract %7 %11 0 -%26 = OpIMul %7 %24 %25 -%27 = OpIAdd %7 %23 %26 -%28 = OpCompositeExtract %7 %11 1 -%29 = OpCompositeExtract %7 %11 1 -%30 = OpIMul %7 %28 %29 -%22 = OpIAdd %7 %27 %30 -%31 = OpCopyObject %9 %8 -%32 = OpExtInst %9 %1 FindUMsb %31 +%23 = OpCompositeConstruct %16 %5 %5 %5 %5 +%24 = OpCompositeConstruct %16 %3 %3 %3 %3 +%22 = OpExtInst %16 %1 FClamp %17 %23 %24 +%27 = OpCompositeExtract %7 %11 0 +%28 = OpCompositeExtract %7 %11 0 +%29 = OpIMul %7 %27 %28 +%30 = OpIAdd %7 %26 %29 +%31 = OpCompositeExtract %7 %11 1 +%32 = OpCompositeExtract %7 %11 1 +%33 = OpIMul %7 %31 %32 +%25 = OpIAdd %7 %30 %33 +%34 = OpCopyObject %9 %8 +%35 = OpExtInst %9 %1 FindUMsb %34 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 9543168794..2643879ee3 100644 --- a/tests/out/wgsl/math-functions.wgsl +++ b/tests/out/wgsl/math-functions.wgsl @@ -5,6 +5,7 @@ fn main() { let b = radians(1.0); let c = degrees(v); let d = radians(v); + let e = saturate(v); let const_dot = dot(vec2(0, 0), vec2(0, 0)); let first_leading_bit_abs = firstLeadingBit(abs(0u)); }