nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-11-23 07:14:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add partial derivative builtins (#2277)

Browse Source 
* Add partial derivative builtins

* [dot] emit derivative control

* Fix fmt

---------

Co-authored-by: Teodor Tanasoaia <28601907+teoxoy@users.noreply.github.com>
This commit is contained in:
Evan Mark Hopkins 2023-03-16 19:15:44 -04:00 committed by GitHub
parent 6fcd04ec06
commit 6db8da72ca
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
21 changed files with 377 additions and 117 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
src/back/dot/mod.rs
View File

@ -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);

36
src/back/glsl/mod.rs
View File

@ -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, ")")?
}

46
src/back/hlsl/writer.rs
View File

@ -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;

9
src/back/msl/writer.rs
View File

@ -1579,11 +1579,12 @@ impl<W: Write> Writer<W> {
}
_ => 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)?;

28
src/back/spv/block.rs
View File

@ -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

19
src/back/wgsl/writer.rs
View File

@ -1590,13 +1590,18 @@ impl<W: Write> Writer<W> {
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)?;

30
src/front/glsl/builtins.rs
View File

@ -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<VectorSize>),
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(),

80
src/front/spv/mod.rs
View File

@ -1129,7 +1129,7 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
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<I: Iterator<Item = u32>> Frontend<I> {
let expr = crate::Expression::Derivative {
axis,
ctrl,
expr: arg_handle,
};
self.lookup_expression.insert(
@ -1294,8 +1295,15 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
($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<I: Iterator<Item = u32>> Frontend<I> {
});
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)?;

4
src/front/wgsl/lower/mod.rs
View File

@ -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);

15
src/front/wgsl/parse/conv.rs
View File

@ -113,11 +113,18 @@ pub fn get_scalar_type(word: &str) -> Option<(crate::ScalarKind, crate::Bytes)>
}
}
pub fn map_derivative_axis(word: &str) -> Option<crate::DerivativeAxis> {
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,
}
}

12
src/lib.rs
View File

@ -977,6 +977,17 @@ pub enum AtomicFunction {
Exchange { compare: Option<Handle<Expression>> },
}
/// 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<Expression>,
},

2
src/proc/typifier.rs
View File

@ -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 {

1
src/valid/analyzer.rs
View File

@ -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(),

2
src/valid/expression.rs
View File

@ -832,7 +832,7 @@ impl super::Validator {
}
ShaderStages::all()
}
E::Derivative { axis: _, expr } => {
E::Derivative { expr, .. } => {
match resolver[expr] {
Ti::Scalar {
kind: Sk::Float, ..

5
src/valid/handles.rs
View File

@ -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 } => {

15
tests/in/standard.wgsl
View File

@ -2,8 +2,17 @@
@fragment
fn derivatives(@builtin(position) foo: vec4<f32>) -> @location(0) vec4<f32> {
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;
}

29
tests/out/glsl/standard.derivatives.Fragment.glsl
View File

@ -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;
}

30
tests/out/hlsl/standard.hlsl
View File

@ -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);
}

29
tests/out/msl/standard.msl
View File

@ -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 };
}

68
tests/out/spv/standard.spvasm
View File

@ -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

30
tests/out/wgsl/standard.wgsl
View File

@ -1,7 +1,29 @@
@fragment
fn derivatives(@builtin(position) foo: vec4<f32>) -> @location(0) vec4<f32> {
let x = dpdx(foo);
let y = dpdy(foo);
let z = fwidth(foo);
return ((x + y) * z);
var x: vec4<f32>;
var y: vec4<f32>;
var z: vec4<f32>;
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);
}