From 6db8da72ca5063612face3b5f79805392b68c9df Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins <85699459+evahop@users.noreply.github.com> Date: Thu, 16 Mar 2023 19:15:44 -0400 Subject: [PATCH] Add partial derivative builtins (#2277) * Add partial derivative builtins * [dot] emit derivative control * Fix fmt --------- Co-authored-by: Teodor Tanasoaia <28601907+teoxoy@users.noreply.github.com> --- src/back/dot/mod.rs | 4 +- src/back/glsl/mod.rs | 36 ++++++--- src/back/hlsl/writer.rs | 46 ++++++++--- src/back/msl/writer.rs | 9 ++- src/back/spv/block.rs | 28 +++++-- src/back/wgsl/writer.rs | 19 +++-- src/front/glsl/builtins.rs | 30 +++---- src/front/spv/mod.rs | 80 ++++++++++++++++--- src/front/wgsl/lower/mod.rs | 4 +- src/front/wgsl/parse/conv.rs | 15 +++- src/lib.rs | 12 +++ src/proc/typifier.rs | 2 +- src/valid/analyzer.rs | 1 + src/valid/expression.rs | 2 +- src/valid/handles.rs | 5 +- tests/in/standard.wgsl | 15 +++- .../glsl/standard.derivatives.Fragment.glsl | 29 ++++++- tests/out/hlsl/standard.hlsl | 30 ++++++- tests/out/msl/standard.msl | 29 ++++++- tests/out/spv/standard.spvasm | 68 +++++++++++----- tests/out/wgsl/standard.wgsl | 30 ++++++- 21 files changed, 377 insertions(+), 117 deletions(-) diff --git a/src/back/dot/mod.rs b/src/back/dot/mod.rs index 4799012ab..1167357e8 100644 --- a/src/back/dot/mod.rs +++ b/src/back/dot/mod.rs @@ -505,9 +505,9 @@ fn write_function_expressions( edges.insert("reject", reject); ("Select".into(), 3) } - E::Derivative { axis, expr } => { + E::Derivative { axis, ctrl, expr } => { edges.insert("", expr); - (format!("d{axis:?}").into(), 8) + (format!("d{axis:?}{ctrl:?}").into(), 8) } E::Relational { fun, argument } => { edges.insert("arg", argument); diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 20c3f7930..44685fb99 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -185,6 +185,10 @@ impl Version { fn supports_integer_functions(&self) -> bool { *self >= Version::Desktop(400) || *self >= Version::new_gles(310) } + + fn supports_derivative_control(&self) -> bool { + *self >= Version::Desktop(450) + } } impl PartialOrd for Version { @@ -2812,18 +2816,28 @@ impl<'a, W: Write> Writer<'a, W> { write!(self.out, ")")? } // `Derivative` is a function call to a glsl provided function - Expression::Derivative { axis, expr } => { - use crate::DerivativeAxis as Da; - - write!( - self.out, - "{}(", - match axis { - Da::X => "dFdx", - Da::Y => "dFdy", - Da::Width => "fwidth", + Expression::Derivative { axis, ctrl, expr } => { + use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl}; + let fun_name = if self.options.version.supports_derivative_control() { + match (axis, ctrl) { + (Axis::X, Ctrl::Coarse) => "dFdxCoarse", + (Axis::X, Ctrl::Fine) => "dFdxFine", + (Axis::X, Ctrl::None) => "dFdx", + (Axis::Y, Ctrl::Coarse) => "dFdyCoarse", + (Axis::Y, Ctrl::Fine) => "dFdyFine", + (Axis::Y, Ctrl::None) => "dFdy", + (Axis::Width, Ctrl::Coarse) => "fwidthCoarse", + (Axis::Width, Ctrl::Fine) => "fwidthFine", + (Axis::Width, Ctrl::None) => "fwidth", } - )?; + } else { + match axis { + Axis::X => "dFdx", + Axis::Y => "dFdy", + Axis::Width => "fwidth", + } + }; + write!(self.out, "{fun_name}(")?; self.write_expr(expr, ctx)?; write!(self.out, ")")? } diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 1f0491c58..d11032bbf 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -133,6 +133,13 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { _ => {} } } + + if let Expression::Derivative { axis, ctrl, expr } = *expr { + use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl}; + if axis == Axis::Width && (ctrl == Ctrl::Coarse || ctrl == Ctrl::Fine) { + self.need_bake_expressions.insert(expr); + } + } } } @@ -2801,17 +2808,34 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { let var_name = &self.names[&NameKey::GlobalVariable(var_handle)]; write!(self.out, "({var_name}) - {offset}) / {stride})")? } - Expression::Derivative { axis, expr } => { - use crate::DerivativeAxis as Da; - - let fun_str = match axis { - Da::X => "ddx", - Da::Y => "ddy", - Da::Width => "fwidth", - }; - write!(self.out, "{fun_str}(")?; - self.write_expr(module, expr, func_ctx)?; - write!(self.out, ")")? + Expression::Derivative { axis, ctrl, expr } => { + use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl}; + if axis == Axis::Width && (ctrl == Ctrl::Coarse || ctrl == Ctrl::Fine) { + let tail = match ctrl { + Ctrl::Coarse => "coarse", + Ctrl::Fine => "fine", + Ctrl::None => unreachable!(), + }; + write!(self.out, "abs(ddx_{tail}(")?; + self.write_expr(module, expr, func_ctx)?; + write!(self.out, ")) + abs(ddy_{tail}(")?; + self.write_expr(module, expr, func_ctx)?; + write!(self.out, "))")? + } else { + let fun_str = match (axis, ctrl) { + (Axis::X, Ctrl::Coarse) => "ddx_coarse", + (Axis::X, Ctrl::Fine) => "ddx_fine", + (Axis::X, Ctrl::None) => "ddx", + (Axis::Y, Ctrl::Coarse) => "ddy_coarse", + (Axis::Y, Ctrl::Fine) => "ddy_fine", + (Axis::Y, Ctrl::None) => "ddy", + (Axis::Width, Ctrl::Coarse | Ctrl::Fine) => unreachable!(), + (Axis::Width, Ctrl::None) => "fwidth", + }; + write!(self.out, "{fun_str}(")?; + self.write_expr(module, expr, func_ctx)?; + write!(self.out, ")")? + } } Expression::Relational { fun, argument } => { use crate::RelationalFunction as Rf; diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index bac178729..ee23ca294 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -1579,11 +1579,12 @@ impl Writer { } _ => return Err(Error::Validation), }, - crate::Expression::Derivative { axis, expr } => { + crate::Expression::Derivative { axis, expr, .. } => { + use crate::DerivativeAxis as Axis; let op = match axis { - crate::DerivativeAxis::X => "dfdx", - crate::DerivativeAxis::Y => "dfdy", - crate::DerivativeAxis::Width => "fwidth", + Axis::X => "dfdx", + Axis::Y => "dfdy", + Axis::Width => "fwidth", }; write!(self.out, "{NAMESPACE}::{op}")?; self.put_call_parameters(iter::once(expr), context)?; diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index ccf597917..c3fa8455e 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -1334,15 +1334,29 @@ impl<'w> BlockContext<'w> { block.body.push(instruction); id } - crate::Expression::Derivative { axis, expr } => { - use crate::DerivativeAxis as Da; - + crate::Expression::Derivative { axis, ctrl, expr } => { + use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl}; + match ctrl { + Ctrl::Coarse | Ctrl::Fine => { + self.writer.require_any( + "DerivativeControl", + &[spirv::Capability::DerivativeControl], + )?; + } + Ctrl::None => {} + } let id = self.gen_id(); let expr_id = self.cached[expr]; - let op = match axis { - Da::X => spirv::Op::DPdx, - Da::Y => spirv::Op::DPdy, - Da::Width => spirv::Op::Fwidth, + let op = match (axis, ctrl) { + (Axis::X, Ctrl::Coarse) => spirv::Op::DPdxCoarse, + (Axis::X, Ctrl::Fine) => spirv::Op::DPdxFine, + (Axis::X, Ctrl::None) => spirv::Op::DPdx, + (Axis::Y, Ctrl::Coarse) => spirv::Op::DPdyCoarse, + (Axis::Y, Ctrl::Fine) => spirv::Op::DPdyFine, + (Axis::Y, Ctrl::None) => spirv::Op::DPdy, + (Axis::Width, Ctrl::Coarse) => spirv::Op::FwidthCoarse, + (Axis::Width, Ctrl::Fine) => spirv::Op::FwidthFine, + (Axis::Width, Ctrl::None) => spirv::Op::Fwidth, }; block .body diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index 7bfb0e964..f24f4a9c2 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -1590,13 +1590,18 @@ impl Writer { self.write_expr(module, condition, func_ctx)?; write!(self.out, ")")? } - Expression::Derivative { axis, expr } => { - use crate::DerivativeAxis as Da; - - let op = match axis { - Da::X => "dpdx", - Da::Y => "dpdy", - Da::Width => "fwidth", + Expression::Derivative { axis, ctrl, expr } => { + use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl}; + let op = match (axis, ctrl) { + (Axis::X, Ctrl::Coarse) => "dpdxCoarse", + (Axis::X, Ctrl::Fine) => "dpdxFine", + (Axis::X, Ctrl::None) => "dpdx", + (Axis::Y, Ctrl::Coarse) => "dpdyCoarse", + (Axis::Y, Ctrl::Fine) => "dpdyFine", + (Axis::Y, Ctrl::None) => "dpdy", + (Axis::Width, Ctrl::Coarse) => "fwidthCoarse", + (Axis::Width, Ctrl::Fine) => "fwidthFine", + (Axis::Width, Ctrl::None) => "fwidth", }; write!(self.out, "{op}(")?; self.write_expr(module, expr, func_ctx)?; diff --git a/src/front/glsl/builtins.rs b/src/front/glsl/builtins.rs index 387ecf012..14e223d6a 100644 --- a/src/front/glsl/builtins.rs +++ b/src/front/glsl/builtins.rs @@ -7,9 +7,10 @@ use super::{ Error, ErrorKind, Frontend, Result, }; use crate::{ - BinaryOperator, Block, Constant, DerivativeAxis, Expression, Handle, ImageClass, - ImageDimension as Dim, ImageQuery, MathFunction, Module, RelationalFunction, SampleLevel, - ScalarKind as Sk, Span, Type, TypeInner, UnaryOperator, VectorSize, + BinaryOperator, Block, Constant, DerivativeAxis as Axis, DerivativeControl as Ctrl, Expression, + Handle, ImageClass, ImageDimension as Dim, ImageQuery, MathFunction, Module, + RelationalFunction, SampleLevel, ScalarKind as Sk, Span, Type, TypeInner, UnaryOperator, + VectorSize, }; impl crate::ScalarKind { @@ -571,15 +572,15 @@ fn inject_standard_builtins( "degrees" => MacroCall::MathFunction(MathFunction::Degrees), "floatBitsToInt" => MacroCall::BitCast(Sk::Sint), "floatBitsToUint" => MacroCall::BitCast(Sk::Uint), - "dFdx" | "dFdxFine" | "dFdxCoarse" => { - MacroCall::Derivate(DerivativeAxis::X) - } - "dFdy" | "dFdyFine" | "dFdyCoarse" => { - MacroCall::Derivate(DerivativeAxis::Y) - } - "fwidth" | "fwidthFine" | "fwidthCoarse" => { - MacroCall::Derivate(DerivativeAxis::Width) - } + "dFdxCoarse" => MacroCall::Derivate(Axis::X, Ctrl::Coarse), + "dFdyCoarse" => MacroCall::Derivate(Axis::Y, Ctrl::Coarse), + "fwidthCoarse" => MacroCall::Derivate(Axis::Width, Ctrl::Coarse), + "dFdxFine" => MacroCall::Derivate(Axis::X, Ctrl::Fine), + "dFdyFine" => MacroCall::Derivate(Axis::Y, Ctrl::Fine), + "fwidthFine" => MacroCall::Derivate(Axis::Width, Ctrl::Fine), + "dFdx" => MacroCall::Derivate(Axis::X, Ctrl::None), + "dFdy" => MacroCall::Derivate(Axis::Y, Ctrl::None), + "fwidth" => MacroCall::Derivate(Axis::Width, Ctrl::None), _ => unreachable!(), }, )) @@ -1661,7 +1662,7 @@ pub enum MacroCall { MixBoolean, Clamp(Option), BitCast(Sk), - Derivate(DerivativeAxis), + Derivate(Axis, Ctrl), Barrier, /// SmoothStep needs a separate variant because it might need it's inputs /// to be splatted depending on the overload @@ -2143,9 +2144,10 @@ impl MacroCall { Span::default(), body, ), - MacroCall::Derivate(axis) => ctx.add_expression( + MacroCall::Derivate(axis, ctrl) => ctx.add_expression( Expression::Derivative { axis, + ctrl, expr: args[0], }, Span::default(), diff --git a/src/front/spv/mod.rs b/src/front/spv/mod.rs index 82a86a36a..ce42be35b 100644 --- a/src/front/spv/mod.rs +++ b/src/front/spv/mod.rs @@ -1129,7 +1129,7 @@ impl> Frontend { block: &mut crate::Block, block_id: spirv::Word, body_idx: usize, - axis: crate::DerivativeAxis, + (axis, ctrl): (crate::DerivativeAxis, crate::DerivativeControl), ) -> Result<(), Error> { let start = self.data_offset; let result_type_id = self.next()?; @@ -1141,6 +1141,7 @@ impl> Frontend { let expr = crate::Expression::Derivative { axis, + ctrl, expr: arg_handle, }; self.lookup_expression.insert( @@ -1294,8 +1295,15 @@ impl> Frontend { ($op:expr, UNARY) => { self.parse_expr_unary_op(ctx, &mut emitter, &mut block, block_id, body_idx, $op) }; - ($axis:expr, DERIVATIVE) => { - self.parse_expr_derivative(ctx, &mut emitter, &mut block, block_id, body_idx, $axis) + ($axis:expr, $ctrl:expr, DERIVATIVE) => { + self.parse_expr_derivative( + ctx, + &mut emitter, + &mut block, + block_id, + body_idx, + ($axis, $ctrl), + ) }; } @@ -3346,14 +3354,68 @@ impl> Frontend { }); body_idx = loop_body_idx; } - Op::DPdx | Op::DPdxFine | Op::DPdxCoarse => { - parse_expr_op!(crate::DerivativeAxis::X, DERIVATIVE)?; + Op::DPdxCoarse => { + parse_expr_op!( + crate::DerivativeAxis::X, + crate::DerivativeControl::Coarse, + DERIVATIVE + )?; } - Op::DPdy | Op::DPdyFine | Op::DPdyCoarse => { - parse_expr_op!(crate::DerivativeAxis::Y, DERIVATIVE)?; + Op::DPdyCoarse => { + parse_expr_op!( + crate::DerivativeAxis::Y, + crate::DerivativeControl::Coarse, + DERIVATIVE + )?; } - Op::Fwidth | Op::FwidthFine | Op::FwidthCoarse => { - parse_expr_op!(crate::DerivativeAxis::Width, DERIVATIVE)?; + Op::FwidthCoarse => { + parse_expr_op!( + crate::DerivativeAxis::Width, + crate::DerivativeControl::Coarse, + DERIVATIVE + )?; + } + Op::DPdxFine => { + parse_expr_op!( + crate::DerivativeAxis::X, + crate::DerivativeControl::Fine, + DERIVATIVE + )?; + } + Op::DPdyFine => { + parse_expr_op!( + crate::DerivativeAxis::Y, + crate::DerivativeControl::Fine, + DERIVATIVE + )?; + } + Op::FwidthFine => { + parse_expr_op!( + crate::DerivativeAxis::Width, + crate::DerivativeControl::Fine, + DERIVATIVE + )?; + } + Op::DPdx => { + parse_expr_op!( + crate::DerivativeAxis::X, + crate::DerivativeControl::None, + DERIVATIVE + )?; + } + Op::DPdy => { + parse_expr_op!( + crate::DerivativeAxis::Y, + crate::DerivativeControl::None, + DERIVATIVE + )?; + } + Op::Fwidth => { + parse_expr_op!( + crate::DerivativeAxis::Width, + crate::DerivativeControl::None, + DERIVATIVE + )?; } Op::ArrayLength => { inst.expect(5)?; diff --git a/src/front/wgsl/lower/mod.rs b/src/front/wgsl/lower/mod.rs index 01a46f28c..f3b157caa 100644 --- a/src/front/wgsl/lower/mod.rs +++ b/src/front/wgsl/lower/mod.rs @@ -1572,12 +1572,12 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { args.finish()?; crate::Expression::Relational { fun, argument } - } else if let Some(axis) = conv::map_derivative_axis(function.name) { + } else if let Some((axis, ctrl)) = conv::map_derivative(function.name) { let mut args = ctx.prepare_args(arguments, 1, span); let expr = self.expression(args.next()?, ctx.reborrow())?; args.finish()?; - crate::Expression::Derivative { axis, expr } + crate::Expression::Derivative { axis, ctrl, expr } } else if let Some(fun) = conv::map_standard_fun(function.name) { let expected = fun.argument_count() as _; let mut args = ctx.prepare_args(arguments, expected, span); diff --git a/src/front/wgsl/parse/conv.rs b/src/front/wgsl/parse/conv.rs index 0e20774dc..a69dbbe47 100644 --- a/src/front/wgsl/parse/conv.rs +++ b/src/front/wgsl/parse/conv.rs @@ -113,11 +113,18 @@ pub fn get_scalar_type(word: &str) -> Option<(crate::ScalarKind, crate::Bytes)> } } -pub fn map_derivative_axis(word: &str) -> Option { +pub fn map_derivative(word: &str) -> Option<(crate::DerivativeAxis, crate::DerivativeControl)> { + use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl}; match word { - "dpdx" => Some(crate::DerivativeAxis::X), - "dpdy" => Some(crate::DerivativeAxis::Y), - "fwidth" => Some(crate::DerivativeAxis::Width), + "dpdxCoarse" => Some((Axis::X, Ctrl::Coarse)), + "dpdyCoarse" => Some((Axis::Y, Ctrl::Coarse)), + "fwidthCoarse" => Some((Axis::Width, Ctrl::Coarse)), + "dpdxFine" => Some((Axis::X, Ctrl::Fine)), + "dpdyFine" => Some((Axis::Y, Ctrl::Fine)), + "fwidthFine" => Some((Axis::Width, Ctrl::Fine)), + "dpdx" => Some((Axis::X, Ctrl::None)), + "dpdy" => Some((Axis::Y, Ctrl::None)), + "fwidth" => Some((Axis::Width, Ctrl::None)), _ => None, } } diff --git a/src/lib.rs b/src/lib.rs index 259a529a4..c1b48b899 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -977,6 +977,17 @@ pub enum AtomicFunction { Exchange { compare: Option> }, } +/// Hint at which precision to compute a derivative. +#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)] +#[cfg_attr(feature = "serialize", derive(Serialize))] +#[cfg_attr(feature = "deserialize", derive(Deserialize))] +#[cfg_attr(feature = "arbitrary", derive(Arbitrary))] +pub enum DerivativeControl { + Coarse, + Fine, + None, +} + /// Axis on which to compute a derivative. #[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)] #[cfg_attr(feature = "serialize", derive(Serialize))] @@ -1388,6 +1399,7 @@ pub enum Expression { /// Compute the derivative on an axis. Derivative { axis: DerivativeAxis, + ctrl: DerivativeControl, //modifier, expr: Handle, }, diff --git a/src/proc/typifier.rs b/src/proc/typifier.rs index 3b9fa1d50..b9ac46831 100644 --- a/src/proc/typifier.rs +++ b/src/proc/typifier.rs @@ -635,7 +635,7 @@ impl<'a> ResolveContext<'a> { }, crate::Expression::AtomicResult { ty, .. } => TypeResolution::Handle(ty), crate::Expression::Select { accept, .. } => past(accept)?.clone(), - crate::Expression::Derivative { axis: _, expr } => past(expr)?.clone(), + crate::Expression::Derivative { expr, .. } => past(expr)?.clone(), crate::Expression::Relational { fun, argument } => match fun { crate::RelationalFunction::All | crate::RelationalFunction::Any => { TypeResolution::Value(Ti::Scalar { diff --git a/src/valid/analyzer.rs b/src/valid/analyzer.rs index 30fb32d07..40d5f95c1 100644 --- a/src/valid/analyzer.rs +++ b/src/valid/analyzer.rs @@ -1015,6 +1015,7 @@ fn uniform_control_flow() { let derivative_expr = expressions.append( E::Derivative { axis: crate::DerivativeAxis::X, + ctrl: crate::DerivativeControl::None, expr: constant_expr, }, Default::default(), diff --git a/src/valid/expression.rs b/src/valid/expression.rs index 01d6910eb..af080fc18 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -832,7 +832,7 @@ impl super::Validator { } ShaderStages::all() } - E::Derivative { axis: _, expr } => { + E::Derivative { expr, .. } => { match resolver[expr] { Ti::Scalar { kind: Sk::Float, .. diff --git a/src/valid/handles.rs b/src/valid/handles.rs index be87e54d4..e3f9fe253 100644 --- a/src/valid/handles.rs +++ b/src/valid/handles.rs @@ -345,10 +345,7 @@ impl super::Validator { .check_dep(accept)? .check_dep(reject)?; } - crate::Expression::Derivative { - axis: _, - expr: argument, - } => { + crate::Expression::Derivative { expr: argument, .. } => { handle.check_dep(argument)?; } crate::Expression::Relational { fun: _, argument } => { diff --git a/tests/in/standard.wgsl b/tests/in/standard.wgsl index 5b40b683a..79f5632ce 100644 --- a/tests/in/standard.wgsl +++ b/tests/in/standard.wgsl @@ -2,8 +2,17 @@ @fragment fn derivatives(@builtin(position) foo: vec4) -> @location(0) vec4 { - let x = dpdx(foo); - let y = dpdy(foo); - let z = fwidth(foo); + var x = dpdxCoarse(foo); + var y = dpdyCoarse(foo); + var z = fwidthCoarse(foo); + + x = dpdxFine(foo); + y = dpdyFine(foo); + z = fwidthFine(foo); + + x = dpdx(foo); + y = dpdy(foo); + z = fwidth(foo); + return (x + y) * z; } diff --git a/tests/out/glsl/standard.derivatives.Fragment.glsl b/tests/out/glsl/standard.derivatives.Fragment.glsl index 4a084add1..331ffb5c0 100644 --- a/tests/out/glsl/standard.derivatives.Fragment.glsl +++ b/tests/out/glsl/standard.derivatives.Fragment.glsl @@ -7,10 +7,31 @@ layout(location = 0) out vec4 _fs2p_location0; void main() { vec4 foo = gl_FragCoord; - vec4 x = dFdx(foo); - vec4 y = dFdy(foo); - vec4 z = fwidth(foo); - _fs2p_location0 = ((x + y) * z); + vec4 x = vec4(0.0); + vec4 y = vec4(0.0); + vec4 z = vec4(0.0); + vec4 _e1 = dFdx(foo); + x = _e1; + vec4 _e3 = dFdy(foo); + y = _e3; + vec4 _e5 = fwidth(foo); + z = _e5; + vec4 _e7 = dFdx(foo); + x = _e7; + vec4 _e8 = dFdy(foo); + y = _e8; + vec4 _e9 = fwidth(foo); + z = _e9; + vec4 _e10 = dFdx(foo); + x = _e10; + vec4 _e11 = dFdy(foo); + y = _e11; + vec4 _e12 = fwidth(foo); + z = _e12; + vec4 _e13 = x; + vec4 _e14 = y; + vec4 _e16 = z; + _fs2p_location0 = ((_e13 + _e14) * _e16); return; } diff --git a/tests/out/hlsl/standard.hlsl b/tests/out/hlsl/standard.hlsl index e4077b5f0..aacb25012 100644 --- a/tests/out/hlsl/standard.hlsl +++ b/tests/out/hlsl/standard.hlsl @@ -6,8 +6,30 @@ struct FragmentInput_derivatives { float4 derivatives(FragmentInput_derivatives fragmentinput_derivatives) : SV_Target0 { float4 foo = fragmentinput_derivatives.foo_1; - float4 x = ddx(foo); - float4 y = ddy(foo); - float4 z = fwidth(foo); - return ((x + y) * z); + float4 x = (float4)0; + float4 y = (float4)0; + float4 z = (float4)0; + + float4 _expr1 = ddx_coarse(foo); + x = _expr1; + float4 _expr3 = ddy_coarse(foo); + y = _expr3; + float4 _expr5 = abs(ddx_coarse(foo)) + abs(ddy_coarse(foo)); + z = _expr5; + float4 _expr7 = ddx_fine(foo); + x = _expr7; + float4 _expr8 = ddy_fine(foo); + y = _expr8; + float4 _expr9 = abs(ddx_fine(foo)) + abs(ddy_fine(foo)); + z = _expr9; + float4 _expr10 = ddx(foo); + x = _expr10; + float4 _expr11 = ddy(foo); + y = _expr11; + float4 _expr12 = fwidth(foo); + z = _expr12; + float4 _expr13 = x; + float4 _expr14 = y; + float4 _expr16 = z; + return ((_expr13 + _expr14) * _expr16); } diff --git a/tests/out/msl/standard.msl b/tests/out/msl/standard.msl index d21930c06..bca5f0cb0 100644 --- a/tests/out/msl/standard.msl +++ b/tests/out/msl/standard.msl @@ -13,8 +13,29 @@ struct derivativesOutput { fragment derivativesOutput derivatives( metal::float4 foo [[position]] ) { - metal::float4 x = metal::dfdx(foo); - metal::float4 y = metal::dfdy(foo); - metal::float4 z = metal::fwidth(foo); - return derivativesOutput { (x + y) * z }; + metal::float4 x = {}; + metal::float4 y = {}; + metal::float4 z = {}; + metal::float4 _e1 = metal::dfdx(foo); + x = _e1; + metal::float4 _e3 = metal::dfdy(foo); + y = _e3; + metal::float4 _e5 = metal::fwidth(foo); + z = _e5; + metal::float4 _e7 = metal::dfdx(foo); + x = _e7; + metal::float4 _e8 = metal::dfdy(foo); + y = _e8; + metal::float4 _e9 = metal::fwidth(foo); + z = _e9; + metal::float4 _e10 = metal::dfdx(foo); + x = _e10; + metal::float4 _e11 = metal::dfdy(foo); + y = _e11; + metal::float4 _e12 = metal::fwidth(foo); + z = _e12; + metal::float4 _e13 = x; + metal::float4 _e14 = y; + metal::float4 _e16 = z; + return derivativesOutput { (_e13 + _e14) * _e16 }; } diff --git a/tests/out/spv/standard.spvasm b/tests/out/spv/standard.spvasm index ef0bac6c9..518be0a22 100644 --- a/tests/out/spv/standard.spvasm +++ b/tests/out/spv/standard.spvasm @@ -1,32 +1,58 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 19 +; Bound: 35 OpCapability Shader +OpCapability DerivativeControl %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Fragment %11 "derivatives" %6 %9 -OpExecutionMode %11 OriginUpperLeft -OpDecorate %6 BuiltIn FragCoord -OpDecorate %9 Location 0 +OpEntryPoint Fragment %18 "derivatives" %13 %16 +OpExecutionMode %18 OriginUpperLeft +OpDecorate %13 BuiltIn FragCoord +OpDecorate %16 Location 0 %2 = OpTypeVoid %4 = OpTypeFloat 32 %3 = OpTypeVector %4 4 -%7 = OpTypePointer Input %3 -%6 = OpVariable %7 Input -%10 = OpTypePointer Output %3 -%9 = OpVariable %10 Output -%12 = OpTypeFunction %2 -%11 = OpFunction %2 None %12 -%5 = OpLabel -%8 = OpLoad %3 %6 -OpBranch %13 -%13 = OpLabel -%14 = OpDPdx %3 %8 -%15 = OpDPdy %3 %8 -%16 = OpFwidth %3 %8 -%17 = OpFAdd %3 %14 %15 -%18 = OpFMul %3 %17 %16 -OpStore %9 %18 +%6 = OpTypePointer Function %3 +%7 = OpConstantNull %3 +%9 = OpConstantNull %3 +%11 = OpConstantNull %3 +%14 = OpTypePointer Input %3 +%13 = OpVariable %14 Input +%17 = OpTypePointer Output %3 +%16 = OpVariable %17 Output +%19 = OpTypeFunction %2 +%18 = OpFunction %2 None %19 +%12 = OpLabel +%5 = OpVariable %6 Function %7 +%8 = OpVariable %6 Function %9 +%10 = OpVariable %6 Function %11 +%15 = OpLoad %3 %13 +OpBranch %20 +%20 = OpLabel +%21 = OpDPdxCoarse %3 %15 +OpStore %5 %21 +%22 = OpDPdyCoarse %3 %15 +OpStore %8 %22 +%23 = OpFwidthCoarse %3 %15 +OpStore %10 %23 +%24 = OpDPdxFine %3 %15 +OpStore %5 %24 +%25 = OpDPdyFine %3 %15 +OpStore %8 %25 +%26 = OpFwidthFine %3 %15 +OpStore %10 %26 +%27 = OpDPdx %3 %15 +OpStore %5 %27 +%28 = OpDPdy %3 %15 +OpStore %8 %28 +%29 = OpFwidth %3 %15 +OpStore %10 %29 +%30 = OpLoad %3 %5 +%31 = OpLoad %3 %8 +%32 = OpFAdd %3 %30 %31 +%33 = OpLoad %3 %10 +%34 = OpFMul %3 %32 %33 +OpStore %16 %34 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/standard.wgsl b/tests/out/wgsl/standard.wgsl index a121d6f9a..80e8f2498 100644 --- a/tests/out/wgsl/standard.wgsl +++ b/tests/out/wgsl/standard.wgsl @@ -1,7 +1,29 @@ @fragment fn derivatives(@builtin(position) foo: vec4) -> @location(0) vec4 { - let x = dpdx(foo); - let y = dpdy(foo); - let z = fwidth(foo); - return ((x + y) * z); + var x: vec4; + var y: vec4; + var z: vec4; + + let _e1 = dpdxCoarse(foo); + x = _e1; + let _e3 = dpdyCoarse(foo); + y = _e3; + let _e5 = fwidthCoarse(foo); + z = _e5; + let _e7 = dpdxFine(foo); + x = _e7; + let _e8 = dpdyFine(foo); + y = _e8; + let _e9 = fwidthFine(foo); + z = _e9; + let _e10 = dpdx(foo); + x = _e10; + let _e11 = dpdy(foo); + y = _e11; + let _e12 = fwidth(foo); + z = _e12; + let _e13 = x; + let _e14 = y; + let _e16 = z; + return ((_e13 + _e14) * _e16); }