diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index dbcda3ada5..fa372ae638 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 { @@ -2829,17 +2833,33 @@ impl<'a, W: Write> Writer<'a, W> { } // `Derivative` is a function call to a glsl provided function Expression::Derivative { axis, expr } => { - use crate::DerivativeAxis as Da; - - write!( - self.out, - "{}(", + use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl}; + let fun_name = if self.options.version.supports_derivative_control() { + match axis { + Axis::X(ctrl) => match ctrl { + Ctrl::Coarse => "dFdxCoarse", + Ctrl::Fine => "dFdxFine", + _ => "dFdx", + }, + Axis::Y(ctrl) => match ctrl { + Ctrl::Coarse => "dFdyCoarse", + Ctrl::Fine => "dFdyFine", + _ => "dFdy", + }, + Axis::Width(ctrl) => match ctrl { + Ctrl::Coarse => "fwidthCoarse", + Ctrl::Fine => "fwidthFine", + _ => "fwidth", + }, + } + } else { match axis { - Da::X => "dFdx", - Da::Y => "dFdy", - Da::Width => "fwidth", + Axis::X(_) => "dFdx", + Axis::Y(_) => "dFdy", + _ => "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 3326c1c96d..465416627d 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -2802,12 +2802,19 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { write!(self.out, "({var_name}) - {offset}) / {stride})")? } Expression::Derivative { axis, expr } => { - use crate::DerivativeAxis as Da; - + use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl}; let fun_str = match axis { - Da::X => "ddx", - Da::Y => "ddy", - Da::Width => "fwidth", + Axis::X(ctrl) => match ctrl { + Ctrl::Coarse => "ddx_coarse", + Ctrl::Fine => "ddx_fine", + _ => "ddx", + }, + Axis::Y(ctrl) => match ctrl { + Ctrl::Coarse => "ddy_coarse", + Ctrl::Fine => "ddy_fine", + _ => "ddy", + }, + _ => "fwidth", }; write!(self.out, "{fun_str}(")?; self.write_expr(module, expr, func_ctx)?; diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index fb4b0929cb..0d32fc5ffd 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -1580,10 +1580,11 @@ impl Writer { _ => return Err(Error::Validation), }, 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", + _ => "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 ccf5979176..7be7935174 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -1335,14 +1335,36 @@ impl<'w> BlockContext<'w> { id } crate::Expression::Derivative { axis, expr } => { - use crate::DerivativeAxis as Da; - + use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl}; + match axis { + Axis::X(ctrl) | Axis::Y(ctrl) | Axis::Width(ctrl) => match ctrl { + Ctrl::Coarse | Ctrl::Fine => { + self.writer.require_any( + "DerivativeControl", + &[spirv::Capability::DerivativeControl], + )?; + } + _ => {} + }, + } 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, + Axis::X(ctrl) => match ctrl { + Ctrl::Coarse => spirv::Op::DPdxCoarse, + Ctrl::Fine => spirv::Op::DPdxFine, + _ => spirv::Op::DPdx, + }, + Axis::Y(ctrl) => match ctrl { + Ctrl::Coarse => spirv::Op::DPdyCoarse, + Ctrl::Fine => spirv::Op::DPdyFine, + _ => spirv::Op::DPdy, + }, + Axis::Width(ctrl) => match ctrl { + Ctrl::Coarse => spirv::Op::FwidthCoarse, + Ctrl::Fine => spirv::Op::FwidthFine, + _ => spirv::Op::Fwidth, + }, }; block .body diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index 828cd79b0d..33dd46cc73 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -1658,12 +1658,23 @@ impl Writer { write!(self.out, ")")? } Expression::Derivative { axis, expr } => { - use crate::DerivativeAxis as Da; - + use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl}; let op = match axis { - Da::X => "dpdx", - Da::Y => "dpdy", - Da::Width => "fwidth", + Axis::X(ctrl) => match ctrl { + Ctrl::Coarse => "dpdxCoarse", + Ctrl::Fine => "dpdxFine", + _ => "dpdx", + }, + Axis::Y(ctrl) => match ctrl { + Ctrl::Coarse => "dpdyCoarse", + Ctrl::Fine => "dpdyFine", + _ => "dpdy", + }, + Axis::Width(ctrl) => match ctrl { + Ctrl::Coarse => "fwidthCoarse", + Ctrl::Fine => "fwidthFine", + _ => "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 ef98e47ea1..12971165a9 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), Barrier, /// SmoothStep needs a separate variant because it might need it's inputs /// to be splatted depending on the overload diff --git a/src/front/glsl/functions.rs b/src/front/glsl/functions.rs index 1f24d404d2..c565b3464a 100644 --- a/src/front/glsl/functions.rs +++ b/src/front/glsl/functions.rs @@ -301,7 +301,6 @@ impl Frontend { Expression::Compose { ty: vector_ty, components: (0..rows as u32) - .into_iter() .map(|r| match r == i { true => value, false => zero, @@ -376,7 +375,6 @@ impl Frontend { components.push(match ori_rows.cmp(&rows) { Ordering::Less => { let components = (0..rows as u32) - .into_iter() .map(|r| { if r < ori_rows as u32 { ctx.add_expression( @@ -415,7 +413,6 @@ impl Frontend { inner: ConstantInner::Composite { ty: vector_ty, components: (0..rows as u32) - .into_iter() .map(|r| match r == i { true => one_constant, false => zero_constant, diff --git a/src/front/spv/mod.rs b/src/front/spv/mod.rs index e5e9d4700f..0c94082cff 100644 --- a/src/front/spv/mod.rs +++ b/src/front/spv/mod.rs @@ -3337,14 +3337,59 @@ 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::DPdyCoarse => { + parse_expr_op!( + crate::DerivativeAxis::Y(crate::DerivativeControl::Coarse), + 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::DPdy | Op::DPdyFine | Op::DPdyCoarse => { - parse_expr_op!(crate::DerivativeAxis::Y, DERIVATIVE)?; + Op::DPdyFine => { + parse_expr_op!( + crate::DerivativeAxis::Y(crate::DerivativeControl::Fine), + DERIVATIVE + )?; } - Op::Fwidth | Op::FwidthFine | Op::FwidthCoarse => { - parse_expr_op!(crate::DerivativeAxis::Width, 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/parse/conv.rs b/src/front/wgsl/parse/conv.rs index 0e20774dc5..26616218b8 100644 --- a/src/front/wgsl/parse/conv.rs +++ b/src/front/wgsl/parse/conv.rs @@ -114,10 +114,17 @@ pub fn get_scalar_type(word: &str) -> Option<(crate::ScalarKind, crate::Bytes)> } pub fn map_derivative_axis(word: &str) -> Option { + 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 e370b81a6f..22a8311658 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -976,15 +976,26 @@ 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))] #[cfg_attr(feature = "deserialize", derive(Deserialize))] #[cfg_attr(feature = "arbitrary", derive(Arbitrary))] pub enum DerivativeAxis { - X, - Y, - Width, + X(DerivativeControl), + Y(DerivativeControl), + Width(DerivativeControl), } /// Built-in shader function for testing relation between values. diff --git a/src/proc/typifier.rs b/src/proc/typifier.rs index 3b9fa1d50b..b9ac468313 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 30fb32d076..9f8653fcf6 100644 --- a/src/valid/analyzer.rs +++ b/src/valid/analyzer.rs @@ -1014,7 +1014,7 @@ fn uniform_control_flow() { // checks the non-uniform control flow let derivative_expr = expressions.append( E::Derivative { - axis: crate::DerivativeAxis::X, + axis: crate::DerivativeAxis::X(crate::DerivativeControl::None), expr: constant_expr, }, Default::default(), diff --git a/src/valid/expression.rs b/src/valid/expression.rs index 01d6910eba..af080fc183 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 be87e54d48..e3f9fe2531 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 5b40b683a3..79f5632cea 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 4a084add13..331ffb5c01 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 e4077b5f04..5c49a82b84 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 = fwidth(foo); + z = _expr5; + float4 _expr7 = ddx_fine(foo); + x = _expr7; + float4 _expr8 = ddy_fine(foo); + y = _expr8; + float4 _expr9 = fwidth(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 d21930c061..bca5f0cb00 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 ef0bac6c90..518be0a22a 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 a121d6f9a0..80e8f24989 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); }