mirror of
https://github.com/gfx-rs/wgpu.git
synced 2025-12-08 21:26:17 +00:00
[naga const-eval] LiteralVector and some demo builtins (#8223)
* [naga] Add `LiteralVector` and `match_literal_vector!` Co-authored-by: SelkoSays <70065171+SelkoSays@users.noreply.github.com> Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com> * [naga] Implement builtins dot, length, distance, normalize in const using LiteralVector Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com> * Update snapshots Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com> --------- Signed-off-by: sagudev <16504129+sagudev@users.noreply.github.com> Co-authored-by: SelkoSays <70065171+SelkoSays@users.noreply.github.com>
This commit is contained in:
Sam
GitHub
parent
42a5760a6c
commit
6e76a98c99
@ -268,6 +268,392 @@ gen_component_wise_extractor! {
|
||||
],
|
||||
}
|
||||
|
||||
/// Vectors with a concrete element type.
|
||||
#[derive(Debug)]
|
||||
enum LiteralVector {
|
||||
F64(ArrayVec<f64, { crate::VectorSize::MAX }>),
|
||||
F32(ArrayVec<f32, { crate::VectorSize::MAX }>),
|
||||
F16(ArrayVec<f16, { crate::VectorSize::MAX }>),
|
||||
U32(ArrayVec<u32, { crate::VectorSize::MAX }>),
|
||||
I32(ArrayVec<i32, { crate::VectorSize::MAX }>),
|
||||
U64(ArrayVec<u64, { crate::VectorSize::MAX }>),
|
||||
I64(ArrayVec<i64, { crate::VectorSize::MAX }>),
|
||||
Bool(ArrayVec<bool, { crate::VectorSize::MAX }>),
|
||||
AbstractInt(ArrayVec<i64, { crate::VectorSize::MAX }>),
|
||||
AbstractFloat(ArrayVec<f64, { crate::VectorSize::MAX }>),
|
||||
}
|
||||
|
||||
impl LiteralVector {
|
||||
#[allow(clippy::missing_const_for_fn, reason = "MSRV")]
|
||||
fn len(&self) -> usize {
|
||||
match *self {
|
||||
LiteralVector::F64(ref v) => v.len(),
|
||||
LiteralVector::F32(ref v) => v.len(),
|
||||
LiteralVector::F16(ref v) => v.len(),
|
||||
LiteralVector::U32(ref v) => v.len(),
|
||||
LiteralVector::I32(ref v) => v.len(),
|
||||
LiteralVector::U64(ref v) => v.len(),
|
||||
LiteralVector::I64(ref v) => v.len(),
|
||||
LiteralVector::Bool(ref v) => v.len(),
|
||||
LiteralVector::AbstractInt(ref v) => v.len(),
|
||||
LiteralVector::AbstractFloat(ref v) => v.len(),
|
||||
}
|
||||
}
|
||||
|
||||
/// Creates [`LiteralVector`] of size 1 from single [`Literal`]
|
||||
fn from_literal(literal: Literal) -> Self {
|
||||
match literal {
|
||||
Literal::F64(e) => Self::F64(ArrayVec::from_iter(iter::once(e))),
|
||||
Literal::F32(e) => Self::F32(ArrayVec::from_iter(iter::once(e))),
|
||||
Literal::U32(e) => Self::U32(ArrayVec::from_iter(iter::once(e))),
|
||||
Literal::I32(e) => Self::I32(ArrayVec::from_iter(iter::once(e))),
|
||||
Literal::U64(e) => Self::U64(ArrayVec::from_iter(iter::once(e))),
|
||||
Literal::I64(e) => Self::I64(ArrayVec::from_iter(iter::once(e))),
|
||||
Literal::Bool(e) => Self::Bool(ArrayVec::from_iter(iter::once(e))),
|
||||
Literal::AbstractInt(e) => Self::AbstractInt(ArrayVec::from_iter(iter::once(e))),
|
||||
Literal::AbstractFloat(e) => Self::AbstractFloat(ArrayVec::from_iter(iter::once(e))),
|
||||
Literal::F16(e) => Self::F16(ArrayVec::from_iter(iter::once(e))),
|
||||
}
|
||||
}
|
||||
|
||||
/// Creates [`LiteralVector`] from [`ArrayVec`] of [`Literal`]s.
|
||||
/// Returns error if components types do not match.
|
||||
/// # Panics
|
||||
/// Panics if vector is empty
|
||||
fn from_literal_vec(
|
||||
components: ArrayVec<Literal, { crate::VectorSize::MAX }>,
|
||||
) -> Result<Self, ConstantEvaluatorError> {
|
||||
assert!(!components.is_empty());
|
||||
Ok(match components[0] {
|
||||
Literal::I32(_) => Self::I32(
|
||||
components
|
||||
.iter()
|
||||
.map(|l| match l {
|
||||
&Literal::I32(v) => Ok(v),
|
||||
// TODO: should we handle abstract int here?
|
||||
_ => Err(ConstantEvaluatorError::InvalidMathArg),
|
||||
})
|
||||
.collect::<Result<_, _>>()?,
|
||||
),
|
||||
Literal::U32(_) => Self::U32(
|
||||
components
|
||||
.iter()
|
||||
.map(|l| match l {
|
||||
&Literal::U32(v) => Ok(v),
|
||||
_ => Err(ConstantEvaluatorError::InvalidMathArg),
|
||||
})
|
||||
.collect::<Result<_, _>>()?,
|
||||
),
|
||||
Literal::I64(_) => Self::I64(
|
||||
components
|
||||
.iter()
|
||||
.map(|l| match l {
|
||||
&Literal::I64(v) => Ok(v),
|
||||
_ => Err(ConstantEvaluatorError::InvalidMathArg),
|
||||
})
|
||||
.collect::<Result<_, _>>()?,
|
||||
),
|
||||
Literal::U64(_) => Self::U64(
|
||||
components
|
||||
.iter()
|
||||
.map(|l| match l {
|
||||
&Literal::U64(v) => Ok(v),
|
||||
_ => Err(ConstantEvaluatorError::InvalidMathArg),
|
||||
})
|
||||
.collect::<Result<_, _>>()?,
|
||||
),
|
||||
Literal::F32(_) => Self::F32(
|
||||
components
|
||||
.iter()
|
||||
.map(|l| match l {
|
||||
&Literal::F32(v) => Ok(v),
|
||||
_ => Err(ConstantEvaluatorError::InvalidMathArg),
|
||||
})
|
||||
.collect::<Result<_, _>>()?,
|
||||
),
|
||||
Literal::F64(_) => Self::F64(
|
||||
components
|
||||
.iter()
|
||||
.map(|l| match l {
|
||||
&Literal::F64(v) => Ok(v),
|
||||
_ => Err(ConstantEvaluatorError::InvalidMathArg),
|
||||
})
|
||||
.collect::<Result<_, _>>()?,
|
||||
),
|
||||
Literal::Bool(_) => Self::Bool(
|
||||
components
|
||||
.iter()
|
||||
.map(|l| match l {
|
||||
&Literal::Bool(v) => Ok(v),
|
||||
_ => Err(ConstantEvaluatorError::InvalidMathArg),
|
||||
})
|
||||
.collect::<Result<_, _>>()?,
|
||||
),
|
||||
Literal::AbstractInt(_) => Self::AbstractInt(
|
||||
components
|
||||
.iter()
|
||||
.map(|l| match l {
|
||||
&Literal::AbstractInt(v) => Ok(v),
|
||||
_ => Err(ConstantEvaluatorError::InvalidMathArg),
|
||||
})
|
||||
.collect::<Result<_, _>>()?,
|
||||
),
|
||||
Literal::AbstractFloat(_) => Self::AbstractFloat(
|
||||
components
|
||||
.iter()
|
||||
.map(|l| match l {
|
||||
&Literal::AbstractFloat(v) => Ok(v),
|
||||
_ => Err(ConstantEvaluatorError::InvalidMathArg),
|
||||
})
|
||||
.collect::<Result<_, _>>()?,
|
||||
),
|
||||
Literal::F16(_) => Self::F16(
|
||||
components
|
||||
.iter()
|
||||
.map(|l| match l {
|
||||
&Literal::F16(v) => Ok(v),
|
||||
_ => Err(ConstantEvaluatorError::InvalidMathArg),
|
||||
})
|
||||
.collect::<Result<_, _>>()?,
|
||||
),
|
||||
})
|
||||
}
|
||||
|
||||
#[allow(dead_code)]
|
||||
/// Returns [`ArrayVec`] of [`Literal`]s
|
||||
fn to_literal_vec(&self) -> ArrayVec<Literal, { crate::VectorSize::MAX }> {
|
||||
match *self {
|
||||
LiteralVector::F64(ref v) => v.iter().map(|e| (Literal::F64(*e))).collect(),
|
||||
LiteralVector::F32(ref v) => v.iter().map(|e| (Literal::F32(*e))).collect(),
|
||||
LiteralVector::F16(ref v) => v.iter().map(|e| (Literal::F16(*e))).collect(),
|
||||
LiteralVector::U32(ref v) => v.iter().map(|e| (Literal::U32(*e))).collect(),
|
||||
LiteralVector::I32(ref v) => v.iter().map(|e| (Literal::I32(*e))).collect(),
|
||||
LiteralVector::U64(ref v) => v.iter().map(|e| (Literal::U64(*e))).collect(),
|
||||
LiteralVector::I64(ref v) => v.iter().map(|e| (Literal::I64(*e))).collect(),
|
||||
LiteralVector::Bool(ref v) => v.iter().map(|e| (Literal::Bool(*e))).collect(),
|
||||
LiteralVector::AbstractInt(ref v) => {
|
||||
v.iter().map(|e| (Literal::AbstractInt(*e))).collect()
|
||||
}
|
||||
LiteralVector::AbstractFloat(ref v) => {
|
||||
v.iter().map(|e| (Literal::AbstractFloat(*e))).collect()
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#[allow(dead_code)]
|
||||
/// Puts self into eval's expressions arena and returns handle to it
|
||||
fn register_as_evaluated_expr(
|
||||
&self,
|
||||
eval: &mut ConstantEvaluator<'_>,
|
||||
span: Span,
|
||||
) -> Result<Handle<Expression>, ConstantEvaluatorError> {
|
||||
let lit_vec = self.to_literal_vec();
|
||||
assert!(!lit_vec.is_empty());
|
||||
let expr = if lit_vec.len() == 1 {
|
||||
Expression::Literal(lit_vec[0])
|
||||
} else {
|
||||
Expression::Compose {
|
||||
ty: eval.types.insert(
|
||||
Type {
|
||||
name: None,
|
||||
inner: TypeInner::Vector {
|
||||
size: match lit_vec.len() {
|
||||
2 => crate::VectorSize::Bi,
|
||||
3 => crate::VectorSize::Tri,
|
||||
4 => crate::VectorSize::Quad,
|
||||
_ => unreachable!(),
|
||||
},
|
||||
scalar: lit_vec[0].scalar(),
|
||||
},
|
||||
},
|
||||
Span::UNDEFINED,
|
||||
),
|
||||
components: lit_vec
|
||||
.iter()
|
||||
.map(|&l| eval.register_evaluated_expr(Expression::Literal(l), span))
|
||||
.collect::<Result<_, _>>()?,
|
||||
}
|
||||
};
|
||||
eval.register_evaluated_expr(expr, span)
|
||||
}
|
||||
}
|
||||
|
||||
/// A macro for matching on [`LiteralVector`] variants.
|
||||
///
|
||||
/// `Float` variant expands to `F16`, `F32`, `F64` and `AbstractFloat`.
|
||||
/// `Integer` variant expands to `I32`, `I64`, `U32`, `U64` and `AbstractInt`.
|
||||
///
|
||||
/// For output both [`Literal`] (fold) and [`LiteralVector`] (map) are supported.
|
||||
///
|
||||
/// Example usage:
|
||||
///
|
||||
/// ```rust,ignore
|
||||
/// match_literal_vector!(match v => Literal {
|
||||
/// F16 => |v| {v.sum()},
|
||||
/// Integer => |v| {v.sum()},
|
||||
/// U32 => |v| -> I32 {v.sum()}, // optionally override return type
|
||||
/// })
|
||||
/// ```
|
||||
///
|
||||
/// ```rust,ignore
|
||||
/// match_literal_vector!(match (e1, e2) => LiteralVector {
|
||||
/// F16 => |e1, e2| {e1+e2},
|
||||
/// Integer => |e1, e2| {e1+e2},
|
||||
/// U32 => |e1, e2| -> I32 {e1+e2}, // optionally override return type
|
||||
/// })
|
||||
/// ```
|
||||
macro_rules! match_literal_vector {
|
||||
(match $lit_vec:expr => $out:ident {
|
||||
$(
|
||||
$ty:ident => |$($var:ident),+| $(-> $ret:ident)? { $body:expr }
|
||||
),+
|
||||
$(,)?
|
||||
}) => {
|
||||
match_literal_vector!(@inner_start $lit_vec; $out; [$($ty),+]; [$({ $($var),+ ; $($ret)? ; $body }),+])
|
||||
};
|
||||
|
||||
(@inner_start
|
||||
$lit_vec:expr;
|
||||
$out:ident;
|
||||
[$($ty:ident),+];
|
||||
[$({ $($var:ident),+ ; $($ret:ident)? ; $body:expr }),+]
|
||||
) => {
|
||||
match_literal_vector!(@inner
|
||||
$lit_vec;
|
||||
$out;
|
||||
[$($ty),+];
|
||||
[] <> [$({ $($var),+ ; $($ret)? ; $body }),+]
|
||||
)
|
||||
};
|
||||
|
||||
(@inner
|
||||
$lit_vec:expr;
|
||||
$out:ident;
|
||||
[$ty:ident $(, $ty1:ident)*];
|
||||
[$({$_ty:ident ; $($_var:ident),+ ; $($_ret:ident)? ; $_body:expr}),*] <>
|
||||
[$({ $($var:ident),+ ; $($ret:ident)? ; $body:expr }),+]
|
||||
) => {
|
||||
match_literal_vector!(@inner
|
||||
$ty;
|
||||
$lit_vec;
|
||||
$out;
|
||||
[$($ty1),*];
|
||||
[$({$_ty ; $($_var),+ ; $($_ret)? ; $_body}),*] <>
|
||||
[$({ $($var),+ ; $($ret)? ; $body }),+]
|
||||
)
|
||||
};
|
||||
(@inner
|
||||
Integer;
|
||||
$lit_vec:expr;
|
||||
$out:ident;
|
||||
[$($ty:ident),*];
|
||||
[$({$_ty:ident ; $($_var:ident),+ ; $($_ret:ident)? ; $_body:expr}),*] <>
|
||||
[
|
||||
{ $($var:ident),+ ; $($ret:ident)? ; $body:expr }
|
||||
$(,{ $($var1:ident),+ ; $($ret1:ident)? ; $body1:expr })*
|
||||
]
|
||||
) => {
|
||||
match_literal_vector!(@inner
|
||||
$lit_vec;
|
||||
$out;
|
||||
[U32, I32, U64, I64, AbstractInt $(, $ty)*];
|
||||
[$({$_ty ; $($_var),+ ; $($_ret)? ; $_body}),*] <>
|
||||
[
|
||||
{ $($var),+ ; $($ret)? ; $body }, // U32
|
||||
{ $($var),+ ; $($ret)? ; $body }, // I32
|
||||
{ $($var),+ ; $($ret)? ; $body }, // U64
|
||||
{ $($var),+ ; $($ret)? ; $body }, // I64
|
||||
{ $($var),+ ; $($ret)? ; $body } // AbstractInt
|
||||
$(,{ $($var1),+ ; $($ret1)? ; $body1 })*
|
||||
]
|
||||
)
|
||||
};
|
||||
(@inner
|
||||
Float;
|
||||
$lit_vec:expr;
|
||||
$out:ident;
|
||||
[$($ty:ident),*];
|
||||
[$({$_ty:ident ; $($_var:ident),+ ; $($_ret:ident)? ; $_body:expr}),*] <>
|
||||
[
|
||||
{ $($var:ident),+ ; $($ret:ident)? ; $body:expr }
|
||||
$(,{ $($var1:ident),+ ; $($ret1:ident)? ; $body1:expr })*
|
||||
]
|
||||
) => {
|
||||
match_literal_vector!(@inner
|
||||
$lit_vec;
|
||||
$out;
|
||||
[F16, F32, F64, AbstractFloat $(, $ty)*];
|
||||
[$({$_ty ; $($_var),+ ; $($_ret)? ; $_body}),*] <>
|
||||
[
|
||||
{ $($var),+ ; $($ret)? ; $body }, // F16
|
||||
{ $($var),+ ; $($ret)? ; $body }, // F32
|
||||
{ $($var),+ ; $($ret)? ; $body }, // F64
|
||||
{ $($var),+ ; $($ret)? ; $body } // AbstractFloat
|
||||
$(,{ $($var1),+ ; $($ret1)? ; $body1 })*
|
||||
]
|
||||
)
|
||||
};
|
||||
(@inner
|
||||
$ty:ident;
|
||||
$lit_vec:expr;
|
||||
$out:ident;
|
||||
[$ty1:ident $(,$ty2:ident)*];
|
||||
[$({$_ty:ident ; $($_var:ident),+ ; $($_ret:ident)? ; $_body:expr}),*] <> [
|
||||
{ $($var:ident),+ ; $($ret:ident)? ; $body:expr }
|
||||
$(, { $($var1:ident),+ ; $($ret1:ident)? ; $body1:expr })*
|
||||
]
|
||||
) => {
|
||||
match_literal_vector!(@inner
|
||||
$ty1;
|
||||
$lit_vec;
|
||||
$out;
|
||||
[$($ty2),*];
|
||||
[
|
||||
$({$_ty ; $($_var),+ ; $($_ret)? ; $_body},)*
|
||||
{ $ty; $($var),+ ; $($ret)? ; $body }
|
||||
] <>
|
||||
[$({ $($var1),+ ; $($ret1)? ; $body1 }),*]
|
||||
|
||||
)
|
||||
};
|
||||
(@inner
|
||||
$ty:ident;
|
||||
$lit_vec:expr;
|
||||
$out:ident;
|
||||
[];
|
||||
[$({$_ty:ident ; $($_var:ident),+ ; $($_ret:ident)? ; $_body:expr}),*] <>
|
||||
[{ $($var:ident),+ ; $($ret:ident)? ; $body:expr }]
|
||||
) => {
|
||||
match_literal_vector!(@inner_finish
|
||||
$lit_vec;
|
||||
$out;
|
||||
[
|
||||
$({ $_ty ; $($_var),+ ; $($_ret)? ; $_body },)*
|
||||
{ $ty; $($var),+ ; $($ret)? ; $body }
|
||||
]
|
||||
)
|
||||
};
|
||||
(@inner_finish
|
||||
$lit_vec:expr;
|
||||
$out:ident;
|
||||
[$({$ty:ident ; $($var:ident),+ ; $($ret:ident)? ; $body:expr}),+]
|
||||
) => {
|
||||
match $lit_vec {
|
||||
$(
|
||||
#[allow(unused_parens)]
|
||||
($(LiteralVector::$ty(ref $var)),+) => { Ok(match_literal_vector!(@expand_ret $out; $ty $(; $ret)? ; $body)) }
|
||||
)+
|
||||
_ => Err(ConstantEvaluatorError::InvalidMathArg),
|
||||
}
|
||||
};
|
||||
(@expand_ret $out:ident; $ty:ident; $body:expr) => {
|
||||
$out::$ty($body)
|
||||
};
|
||||
(@expand_ret $out:ident; $_ty:ident; $ret:ident; $body:expr) => {
|
||||
$out::$ret($body)
|
||||
};
|
||||
}
|
||||
|
||||
#[derive(Debug)]
|
||||
enum Behavior<'a> {
|
||||
Wgsl(WgslRestrictions<'a>),
|
||||
@ -1349,16 +1735,105 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
self.packed_dot_product(arg, arg1.unwrap(), span, false)
|
||||
}
|
||||
crate::MathFunction::Cross => self.cross_product(arg, arg1.unwrap(), span),
|
||||
crate::MathFunction::Dot => {
|
||||
// https://www.w3.org/TR/WGSL/#dot-builtin
|
||||
let e1 = self.extract_vec(arg, false)?;
|
||||
let e2 = self.extract_vec(arg1.unwrap(), false)?;
|
||||
if e1.len() != e2.len() {
|
||||
return Err(ConstantEvaluatorError::InvalidMathArg);
|
||||
}
|
||||
|
||||
fn int_dot<P>(a: &[P], b: &[P]) -> Result<P, ConstantEvaluatorError>
|
||||
where
|
||||
P: num_traits::PrimInt + num_traits::CheckedAdd + num_traits::CheckedMul,
|
||||
{
|
||||
a.iter()
|
||||
.zip(b.iter())
|
||||
.map(|(&aa, bb)| aa.checked_mul(bb))
|
||||
.try_fold(P::zero(), |acc, x| {
|
||||
if let Some(x) = x {
|
||||
acc.checked_add(&x)
|
||||
} else {
|
||||
None
|
||||
}
|
||||
})
|
||||
.ok_or(ConstantEvaluatorError::Overflow(
|
||||
"in dot built-in".to_string(),
|
||||
))
|
||||
}
|
||||
|
||||
let result = match_literal_vector!(match (e1, e2) => Literal {
|
||||
Float => |e1, e2| { e1.iter().zip(e2.iter()).map(|(&aa, &bb)| aa * bb).sum() },
|
||||
Integer => |e1, e2| { int_dot(e1, e2)? },
|
||||
})?;
|
||||
self.register_evaluated_expr(Expression::Literal(result), span)
|
||||
}
|
||||
crate::MathFunction::Length => {
|
||||
// https://www.w3.org/TR/WGSL/#length-builtin
|
||||
let e1 = self.extract_vec(arg, true)?;
|
||||
|
||||
fn float_length<F>(e: &[F]) -> F
|
||||
where
|
||||
F: core::ops::Mul<F>,
|
||||
F: num_traits::Float + iter::Sum,
|
||||
{
|
||||
e.iter().map(|&ei| ei * ei).sum::<F>().sqrt()
|
||||
}
|
||||
|
||||
let result = match_literal_vector!(match e1 => Literal {
|
||||
Float => |e1| { float_length(e1) },
|
||||
})?;
|
||||
self.register_evaluated_expr(Expression::Literal(result), span)
|
||||
}
|
||||
crate::MathFunction::Distance => {
|
||||
// https://www.w3.org/TR/WGSL/#distance-builtin
|
||||
let e1 = self.extract_vec(arg, true)?;
|
||||
let e2 = self.extract_vec(arg1.unwrap(), true)?;
|
||||
if e1.len() != e2.len() {
|
||||
return Err(ConstantEvaluatorError::InvalidMathArg);
|
||||
}
|
||||
|
||||
fn float_distance<F>(a: &[F], b: &[F]) -> F
|
||||
where
|
||||
F: core::ops::Mul<F>,
|
||||
F: num_traits::Float + iter::Sum + core::ops::Sub,
|
||||
{
|
||||
a.iter()
|
||||
.zip(b.iter())
|
||||
.map(|(&aa, &bb)| aa - bb)
|
||||
.map(|ei| ei * ei)
|
||||
.sum::<F>()
|
||||
.sqrt()
|
||||
}
|
||||
let result = match_literal_vector!(match (e1, e2) => Literal {
|
||||
Float => |e1, e2| { float_distance(e1, e2) },
|
||||
})?;
|
||||
self.register_evaluated_expr(Expression::Literal(result), span)
|
||||
}
|
||||
crate::MathFunction::Normalize => {
|
||||
// https://www.w3.org/TR/WGSL/#normalize-builtin
|
||||
let e1 = self.extract_vec(arg, true)?;
|
||||
|
||||
fn float_normalize<F>(e: &[F]) -> ArrayVec<F, { crate::VectorSize::MAX }>
|
||||
where
|
||||
F: core::ops::Mul<F>,
|
||||
F: num_traits::Float + iter::Sum,
|
||||
{
|
||||
let len = e.iter().map(|&ei| ei * ei).sum::<F>().sqrt();
|
||||
e.iter().map(|&ei| ei / len).collect()
|
||||
}
|
||||
|
||||
let result = match_literal_vector!(match e1 => LiteralVector {
|
||||
Float => |e1| { float_normalize(e1) },
|
||||
})?;
|
||||
result.register_as_evaluated_expr(self, span)
|
||||
}
|
||||
|
||||
// unimplemented
|
||||
crate::MathFunction::Modf
|
||||
| crate::MathFunction::Frexp
|
||||
| crate::MathFunction::Ldexp
|
||||
| crate::MathFunction::Dot
|
||||
| crate::MathFunction::Outer
|
||||
| crate::MathFunction::Distance
|
||||
| crate::MathFunction::Length
|
||||
| crate::MathFunction::Normalize
|
||||
| crate::MathFunction::FaceForward
|
||||
| crate::MathFunction::Reflect
|
||||
| crate::MathFunction::Refract
|
||||
@ -1434,8 +1909,8 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
) -> Result<Handle<Expression>, ConstantEvaluatorError> {
|
||||
use Literal as Li;
|
||||
|
||||
let (a, ty) = self.extract_vec::<3>(a)?;
|
||||
let (b, _) = self.extract_vec::<3>(b)?;
|
||||
let (a, ty) = self.extract_vec_with_size::<3>(a)?;
|
||||
let (b, _) = self.extract_vec_with_size::<3>(b)?;
|
||||
|
||||
let product = match (a, b) {
|
||||
(
|
||||
@ -1502,7 +1977,7 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
/// values.
|
||||
///
|
||||
/// Also return the type handle from the `Compose` expression.
|
||||
fn extract_vec<const N: usize>(
|
||||
fn extract_vec_with_size<const N: usize>(
|
||||
&mut self,
|
||||
expr: Handle<Expression>,
|
||||
) -> Result<([Literal; N], Handle<Type>), ConstantEvaluatorError> {
|
||||
@ -1526,6 +2001,39 @@ impl<'a> ConstantEvaluator<'a> {
|
||||
Ok((value, ty))
|
||||
}
|
||||
|
||||
/// Extract the values of a `vecN` from `expr`.
|
||||
///
|
||||
/// Return the value of `expr`, whose type is `vecN<S>` for some
|
||||
/// vector size `N` and scalar `S`, as an array of `N` [`Literal`]
|
||||
/// values.
|
||||
///
|
||||
/// Also return the type handle from the `Compose` expression.
|
||||
fn extract_vec(
|
||||
&mut self,
|
||||
expr: Handle<Expression>,
|
||||
allow_single: bool,
|
||||
) -> Result<LiteralVector, ConstantEvaluatorError> {
|
||||
let span = self.expressions.get_span(expr);
|
||||
let expr = self.eval_zero_value_and_splat(expr, span)?;
|
||||
|
||||
match self.expressions[expr] {
|
||||
Expression::Literal(literal) if allow_single => {
|
||||
Ok(LiteralVector::from_literal(literal))
|
||||
}
|
||||
Expression::Compose { ty, ref components } => {
|
||||
let components: ArrayVec<Literal, { crate::VectorSize::MAX }> =
|
||||
crate::proc::flatten_compose(ty, components, self.expressions, self.types)
|
||||
.map(|expr| match self.expressions[expr] {
|
||||
Expression::Literal(l) => Ok(l),
|
||||
_ => Err(ConstantEvaluatorError::InvalidMathArg),
|
||||
})
|
||||
.collect::<Result<_, ConstantEvaluatorError>>()?;
|
||||
LiteralVector::from_literal_vec(components)
|
||||
}
|
||||
_ => Err(ConstantEvaluatorError::InvalidMathArg),
|
||||
}
|
||||
}
|
||||
|
||||
fn array_length(
|
||||
&mut self,
|
||||
array: Handle<Expression>,
|
||||
|
||||
@ -20,10 +20,7 @@ int test_integer_dot_product() {
|
||||
uvec3 a_3_ = uvec3(1u);
|
||||
uvec3 b_3_ = uvec3(1u);
|
||||
uint c_3_ = ( + a_3_.x * b_3_.x + a_3_.y * b_3_.y + a_3_.z * b_3_.z);
|
||||
ivec4 _e11 = ivec4(4);
|
||||
ivec4 _e13 = ivec4(2);
|
||||
int c_4_ = ( + _e11.x * _e13.x + _e11.y * _e13.y + _e11.z * _e13.z + _e11.w * _e13.w);
|
||||
return c_4_;
|
||||
return 32;
|
||||
}
|
||||
|
||||
uint test_packed_integer_dot_product() {
|
||||
|
||||
@ -64,7 +64,6 @@ void main() {
|
||||
vec4 g = refract(v, v, 1.0);
|
||||
ivec4 sign_b = ivec4(-1, -1, -1, -1);
|
||||
vec4 sign_d = vec4(-1.0, -1.0, -1.0, -1.0);
|
||||
int const_dot = ( + ivec2(0).x * ivec2(0).x + ivec2(0).y * ivec2(0).y);
|
||||
ivec2 flb_b = ivec2(-1, -1);
|
||||
uvec2 flb_c = uvec2(0u, 0u);
|
||||
ivec2 ftb_c = ivec2(0, 0);
|
||||
@ -88,12 +87,12 @@ void main() {
|
||||
int frexp_c = naga_frexp(1.5).exp_;
|
||||
int frexp_d = naga_frexp(vec4(1.5, 1.5, 1.5, 1.5)).exp_.x;
|
||||
float quantizeToF16_a = unpackHalf2x16(packHalf2x16(vec2(1.0))).x;
|
||||
vec2 _e120 = vec2(1.0, 1.0);
|
||||
vec2 quantizeToF16_b = unpackHalf2x16(packHalf2x16(_e120));
|
||||
vec3 _e125 = vec3(1.0, 1.0, 1.0);
|
||||
vec3 quantizeToF16_c = vec3(unpackHalf2x16(packHalf2x16(_e125.xy)), unpackHalf2x16(packHalf2x16(_e125.zz)).x);
|
||||
vec4 _e131 = vec4(1.0, 1.0, 1.0, 1.0);
|
||||
vec4 quantizeToF16_d = vec4(unpackHalf2x16(packHalf2x16(_e131.xy)), unpackHalf2x16(packHalf2x16(_e131.zw)));
|
||||
vec2 _e118 = vec2(1.0, 1.0);
|
||||
vec2 quantizeToF16_b = unpackHalf2x16(packHalf2x16(_e118));
|
||||
vec3 _e123 = vec3(1.0, 1.0, 1.0);
|
||||
vec3 quantizeToF16_c = vec3(unpackHalf2x16(packHalf2x16(_e123.xy)), unpackHalf2x16(packHalf2x16(_e123.zz)).x);
|
||||
vec4 _e129 = vec4(1.0, 1.0, 1.0, 1.0);
|
||||
vec4 quantizeToF16_d = vec4(unpackHalf2x16(packHalf2x16(_e129.xy)), unpackHalf2x16(packHalf2x16(_e129.zw)));
|
||||
return;
|
||||
}
|
||||
|
||||
|
||||
@ -14,8 +14,7 @@ int test_integer_dot_product()
|
||||
uint3 a_3_ = (1u).xxx;
|
||||
uint3 b_3_ = (1u).xxx;
|
||||
uint c_3_ = dot(a_3_, b_3_);
|
||||
int c_4_ = dot((int(4)).xxxx, (int(2)).xxxx);
|
||||
return c_4_;
|
||||
return int(32);
|
||||
}
|
||||
|
||||
uint test_packed_integer_dot_product()
|
||||
|
||||
@ -63,10 +63,6 @@ _frexp_result_vec4_f32_ naga_frexp(float4 arg) {
|
||||
return result;
|
||||
}
|
||||
|
||||
int2 ZeroValueint2() {
|
||||
return (int2)0;
|
||||
}
|
||||
|
||||
void main()
|
||||
{
|
||||
float4 v = (0.0).xxxx;
|
||||
@ -78,7 +74,6 @@ void main()
|
||||
float4 g = refract(v, v, 1.0);
|
||||
int4 sign_b = int4(int(-1), int(-1), int(-1), int(-1));
|
||||
float4 sign_d = float4(-1.0, -1.0, -1.0, -1.0);
|
||||
int const_dot = dot(ZeroValueint2(), ZeroValueint2());
|
||||
int2 flb_b = int2(int(-1), int(-1));
|
||||
uint2 flb_c = uint2(0u, 0u);
|
||||
int2 ftb_c = int2(int(0), int(0));
|
||||
|
||||
@ -21,10 +21,7 @@ int test_integer_dot_product(
|
||||
metal::uint3 a_3_ = metal::uint3(1u);
|
||||
metal::uint3 b_3_ = metal::uint3(1u);
|
||||
uint c_3_ = ( + a_3_.x * b_3_.x + a_3_.y * b_3_.y + a_3_.z * b_3_.z);
|
||||
metal::int4 _e11 = metal::int4(4);
|
||||
metal::int4 _e13 = metal::int4(2);
|
||||
int c_4_ = ( + _e11.x * _e13.x + _e11.y * _e13.y + _e11.z * _e13.z + _e11.w * _e13.w);
|
||||
return c_4_;
|
||||
return 32;
|
||||
}
|
||||
|
||||
uint test_packed_integer_dot_product(
|
||||
|
||||
@ -66,7 +66,6 @@ fragment void main_(
|
||||
metal::float4 g = metal::refract(v, v, 1.0);
|
||||
metal::int4 sign_b = metal::int4(-1, -1, -1, -1);
|
||||
metal::float4 sign_d = metal::float4(-1.0, -1.0, -1.0, -1.0);
|
||||
int const_dot = ( + metal::int2 {}.x * metal::int2 {}.x + metal::int2 {}.y * metal::int2 {}.y);
|
||||
metal::int2 flb_b = metal::int2(-1, -1);
|
||||
metal::uint2 flb_c = metal::uint2(0u, 0u);
|
||||
metal::int2 ftb_c = metal::int2(0, 0);
|
||||
|
||||
@ -1,15 +1,15 @@
|
||||
; SPIR-V
|
||||
; Version: 1.1
|
||||
; Generator: rspirv
|
||||
; Bound: 95
|
||||
; Bound: 75
|
||||
OpCapability Shader
|
||||
OpCapability DotProduct
|
||||
OpCapability DotProductInput4x8BitPacked
|
||||
OpExtension "SPV_KHR_integer_dot_product"
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %89 "main"
|
||||
OpExecutionMode %89 LocalSize 1 1 1
|
||||
OpEntryPoint GLCompute %69 "main"
|
||||
OpExecutionMode %69 LocalSize 1 1 1
|
||||
%2 = OpTypeVoid
|
||||
%4 = OpTypeFloat 32
|
||||
%3 = OpTypeVector %4 2
|
||||
@ -27,22 +27,18 @@ OpExecutionMode %89 LocalSize 1 1 1
|
||||
%22 = OpConstant %6 1
|
||||
%23 = OpTypeVector %6 3
|
||||
%24 = OpConstantComposite %23 %22 %22 %22
|
||||
%25 = OpConstant %5 4
|
||||
%26 = OpTypeVector %5 4
|
||||
%27 = OpConstantComposite %26 %25 %25 %25 %25
|
||||
%28 = OpConstant %5 2
|
||||
%29 = OpConstantComposite %26 %28 %28 %28 %28
|
||||
%32 = OpConstantNull %5
|
||||
%41 = OpConstantNull %6
|
||||
%71 = OpTypeFunction %6
|
||||
%72 = OpConstant %6 2
|
||||
%73 = OpConstant %6 3
|
||||
%74 = OpConstant %6 4
|
||||
%75 = OpConstant %6 5
|
||||
%76 = OpConstant %6 6
|
||||
%77 = OpConstant %6 7
|
||||
%78 = OpConstant %6 8
|
||||
%90 = OpTypeFunction %2
|
||||
%25 = OpConstant %5 32
|
||||
%28 = OpConstantNull %5
|
||||
%37 = OpConstantNull %6
|
||||
%51 = OpTypeFunction %6
|
||||
%52 = OpConstant %6 2
|
||||
%53 = OpConstant %6 3
|
||||
%54 = OpConstant %6 4
|
||||
%55 = OpConstant %6 5
|
||||
%56 = OpConstant %6 6
|
||||
%57 = OpConstant %6 7
|
||||
%58 = OpConstant %6 8
|
||||
%70 = OpTypeFunction %2
|
||||
%8 = OpFunction %3 None %9
|
||||
%7 = OpLabel
|
||||
OpBranch %14
|
||||
@ -52,66 +48,50 @@ OpReturnValue %15
|
||||
OpFunctionEnd
|
||||
%17 = OpFunction %5 None %18
|
||||
%16 = OpLabel
|
||||
OpBranch %30
|
||||
%30 = OpLabel
|
||||
%33 = OpCompositeExtract %5 %21 0
|
||||
%34 = OpCompositeExtract %5 %21 0
|
||||
OpBranch %26
|
||||
%26 = OpLabel
|
||||
%29 = OpCompositeExtract %5 %21 0
|
||||
%30 = OpCompositeExtract %5 %21 0
|
||||
%31 = OpIMul %5 %29 %30
|
||||
%32 = OpIAdd %5 %28 %31
|
||||
%33 = OpCompositeExtract %5 %21 1
|
||||
%34 = OpCompositeExtract %5 %21 1
|
||||
%35 = OpIMul %5 %33 %34
|
||||
%36 = OpIAdd %5 %32 %35
|
||||
%37 = OpCompositeExtract %5 %21 1
|
||||
%38 = OpCompositeExtract %5 %21 1
|
||||
%39 = OpIMul %5 %37 %38
|
||||
%31 = OpIAdd %5 %36 %39
|
||||
%42 = OpCompositeExtract %6 %24 0
|
||||
%43 = OpCompositeExtract %6 %24 0
|
||||
%27 = OpIAdd %5 %32 %35
|
||||
%38 = OpCompositeExtract %6 %24 0
|
||||
%39 = OpCompositeExtract %6 %24 0
|
||||
%40 = OpIMul %6 %38 %39
|
||||
%41 = OpIAdd %6 %37 %40
|
||||
%42 = OpCompositeExtract %6 %24 1
|
||||
%43 = OpCompositeExtract %6 %24 1
|
||||
%44 = OpIMul %6 %42 %43
|
||||
%45 = OpIAdd %6 %41 %44
|
||||
%46 = OpCompositeExtract %6 %24 1
|
||||
%47 = OpCompositeExtract %6 %24 1
|
||||
%46 = OpCompositeExtract %6 %24 2
|
||||
%47 = OpCompositeExtract %6 %24 2
|
||||
%48 = OpIMul %6 %46 %47
|
||||
%49 = OpIAdd %6 %45 %48
|
||||
%50 = OpCompositeExtract %6 %24 2
|
||||
%51 = OpCompositeExtract %6 %24 2
|
||||
%52 = OpIMul %6 %50 %51
|
||||
%40 = OpIAdd %6 %49 %52
|
||||
%54 = OpCompositeExtract %5 %27 0
|
||||
%55 = OpCompositeExtract %5 %29 0
|
||||
%56 = OpIMul %5 %54 %55
|
||||
%57 = OpIAdd %5 %32 %56
|
||||
%58 = OpCompositeExtract %5 %27 1
|
||||
%59 = OpCompositeExtract %5 %29 1
|
||||
%60 = OpIMul %5 %58 %59
|
||||
%61 = OpIAdd %5 %57 %60
|
||||
%62 = OpCompositeExtract %5 %27 2
|
||||
%63 = OpCompositeExtract %5 %29 2
|
||||
%64 = OpIMul %5 %62 %63
|
||||
%65 = OpIAdd %5 %61 %64
|
||||
%66 = OpCompositeExtract %5 %27 3
|
||||
%67 = OpCompositeExtract %5 %29 3
|
||||
%68 = OpIMul %5 %66 %67
|
||||
%53 = OpIAdd %5 %65 %68
|
||||
OpReturnValue %53
|
||||
%36 = OpIAdd %6 %45 %48
|
||||
OpReturnValue %25
|
||||
OpFunctionEnd
|
||||
%70 = OpFunction %6 None %71
|
||||
%69 = OpLabel
|
||||
OpBranch %79
|
||||
%79 = OpLabel
|
||||
%80 = OpSDot %5 %22 %72 PackedVectorFormat4x8Bit
|
||||
%81 = OpUDot %6 %73 %74 PackedVectorFormat4x8Bit
|
||||
%82 = OpIAdd %6 %75 %81
|
||||
%83 = OpIAdd %6 %76 %81
|
||||
%84 = OpSDot %5 %82 %83 PackedVectorFormat4x8Bit
|
||||
%85 = OpIAdd %6 %77 %81
|
||||
%86 = OpIAdd %6 %78 %81
|
||||
%87 = OpUDot %6 %85 %86 PackedVectorFormat4x8Bit
|
||||
OpReturnValue %87
|
||||
%50 = OpFunction %6 None %51
|
||||
%49 = OpLabel
|
||||
OpBranch %59
|
||||
%59 = OpLabel
|
||||
%60 = OpSDot %5 %22 %52 PackedVectorFormat4x8Bit
|
||||
%61 = OpUDot %6 %53 %54 PackedVectorFormat4x8Bit
|
||||
%62 = OpIAdd %6 %55 %61
|
||||
%63 = OpIAdd %6 %56 %61
|
||||
%64 = OpSDot %5 %62 %63 PackedVectorFormat4x8Bit
|
||||
%65 = OpIAdd %6 %57 %61
|
||||
%66 = OpIAdd %6 %58 %61
|
||||
%67 = OpUDot %6 %65 %66 PackedVectorFormat4x8Bit
|
||||
OpReturnValue %67
|
||||
OpFunctionEnd
|
||||
%89 = OpFunction %2 None %90
|
||||
%88 = OpLabel
|
||||
OpBranch %91
|
||||
%91 = OpLabel
|
||||
%92 = OpFunctionCall %3 %8
|
||||
%93 = OpFunctionCall %5 %17
|
||||
%94 = OpFunctionCall %6 %70
|
||||
%69 = OpFunction %2 None %70
|
||||
%68 = OpLabel
|
||||
OpBranch %71
|
||||
%71 = OpLabel
|
||||
%72 = OpFunctionCall %3 %8
|
||||
%73 = OpFunctionCall %5 %17
|
||||
%74 = OpFunctionCall %6 %50
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -1,7 +1,7 @@
|
||||
; SPIR-V
|
||||
; Version: 1.1
|
||||
; Generator: rspirv
|
||||
; Bound: 95
|
||||
; Bound: 85
|
||||
OpCapability Shader
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
@ -40,76 +40,66 @@ OpMemberDecorate %15 1 Offset 16
|
||||
%24 = OpConstantComposite %6 %23 %23 %23 %23
|
||||
%25 = OpConstant %3 -1
|
||||
%26 = OpConstantComposite %4 %25 %25 %25 %25
|
||||
%27 = OpConstantNull %7
|
||||
%27 = OpConstant %5 0
|
||||
%28 = OpConstant %9 4294967295
|
||||
%29 = OpConstantComposite %7 %23 %23
|
||||
%30 = OpConstant %9 0
|
||||
%31 = OpConstantComposite %8 %30 %30
|
||||
%32 = OpConstant %5 0
|
||||
%33 = OpConstantComposite %7 %32 %32
|
||||
%34 = OpConstant %9 32
|
||||
%35 = OpConstant %5 32
|
||||
%36 = OpConstantComposite %8 %34 %34
|
||||
%37 = OpConstantComposite %7 %35 %35
|
||||
%38 = OpConstant %9 31
|
||||
%39 = OpConstantComposite %8 %38 %38
|
||||
%40 = OpConstant %5 2
|
||||
%41 = OpConstant %3 2
|
||||
%42 = OpConstantComposite %10 %20 %41
|
||||
%43 = OpConstant %5 3
|
||||
%44 = OpConstant %5 4
|
||||
%45 = OpConstantComposite %7 %43 %44
|
||||
%46 = OpConstant %3 1.5
|
||||
%47 = OpConstantComposite %10 %46 %46
|
||||
%48 = OpConstantComposite %4 %46 %46 %46 %46
|
||||
%49 = OpConstantComposite %10 %20 %20
|
||||
%50 = OpConstantComposite %16 %20 %20 %20
|
||||
%51 = OpConstantComposite %4 %20 %20 %20 %20
|
||||
%58 = OpConstantComposite %4 %20 %20 %20 %20
|
||||
%61 = OpConstantNull %5
|
||||
%32 = OpConstantComposite %7 %27 %27
|
||||
%33 = OpConstant %9 32
|
||||
%34 = OpConstant %5 32
|
||||
%35 = OpConstantComposite %8 %33 %33
|
||||
%36 = OpConstantComposite %7 %34 %34
|
||||
%37 = OpConstant %9 31
|
||||
%38 = OpConstantComposite %8 %37 %37
|
||||
%39 = OpConstant %5 2
|
||||
%40 = OpConstant %3 2
|
||||
%41 = OpConstantComposite %10 %20 %40
|
||||
%42 = OpConstant %5 3
|
||||
%43 = OpConstant %5 4
|
||||
%44 = OpConstantComposite %7 %42 %43
|
||||
%45 = OpConstant %3 1.5
|
||||
%46 = OpConstantComposite %10 %45 %45
|
||||
%47 = OpConstantComposite %4 %45 %45 %45 %45
|
||||
%48 = OpConstantComposite %10 %20 %20
|
||||
%49 = OpConstantComposite %16 %20 %20 %20
|
||||
%50 = OpConstantComposite %4 %20 %20 %20 %20
|
||||
%57 = OpConstantComposite %4 %20 %20 %20 %20
|
||||
%18 = OpFunction %2 None %19
|
||||
%17 = OpLabel
|
||||
OpBranch %52
|
||||
%52 = OpLabel
|
||||
%53 = OpExtInst %3 %1 Degrees %20
|
||||
%54 = OpExtInst %3 %1 Radians %20
|
||||
%55 = OpExtInst %4 %1 Degrees %22
|
||||
%56 = OpExtInst %4 %1 Radians %22
|
||||
%57 = OpExtInst %4 %1 FClamp %22 %22 %58
|
||||
%59 = OpExtInst %4 %1 Refract %22 %22 %20
|
||||
%62 = OpCompositeExtract %5 %27 0
|
||||
%63 = OpCompositeExtract %5 %27 0
|
||||
%64 = OpIMul %5 %62 %63
|
||||
%65 = OpIAdd %5 %61 %64
|
||||
%66 = OpCompositeExtract %5 %27 1
|
||||
%67 = OpCompositeExtract %5 %27 1
|
||||
%68 = OpIMul %5 %66 %67
|
||||
%60 = OpIAdd %5 %65 %68
|
||||
%69 = OpExtInst %3 %1 Ldexp %20 %40
|
||||
%70 = OpExtInst %10 %1 Ldexp %42 %45
|
||||
%71 = OpExtInst %11 %1 ModfStruct %46
|
||||
%72 = OpExtInst %11 %1 ModfStruct %46
|
||||
%73 = OpCompositeExtract %3 %72 0
|
||||
%74 = OpExtInst %11 %1 ModfStruct %46
|
||||
%75 = OpCompositeExtract %3 %74 1
|
||||
%76 = OpExtInst %12 %1 ModfStruct %47
|
||||
%77 = OpExtInst %13 %1 ModfStruct %48
|
||||
%78 = OpCompositeExtract %4 %77 1
|
||||
%79 = OpCompositeExtract %3 %78 0
|
||||
%80 = OpExtInst %12 %1 ModfStruct %47
|
||||
%81 = OpCompositeExtract %10 %80 0
|
||||
%82 = OpCompositeExtract %3 %81 1
|
||||
%83 = OpExtInst %14 %1 FrexpStruct %46
|
||||
%84 = OpExtInst %14 %1 FrexpStruct %46
|
||||
%85 = OpCompositeExtract %3 %84 0
|
||||
%86 = OpExtInst %14 %1 FrexpStruct %46
|
||||
%87 = OpCompositeExtract %5 %86 1
|
||||
%88 = OpExtInst %15 %1 FrexpStruct %48
|
||||
%89 = OpCompositeExtract %6 %88 1
|
||||
%90 = OpCompositeExtract %5 %89 0
|
||||
%91 = OpQuantizeToF16 %3 %20
|
||||
%92 = OpQuantizeToF16 %10 %49
|
||||
%93 = OpQuantizeToF16 %16 %50
|
||||
%94 = OpQuantizeToF16 %4 %51
|
||||
OpBranch %51
|
||||
%51 = OpLabel
|
||||
%52 = OpExtInst %3 %1 Degrees %20
|
||||
%53 = OpExtInst %3 %1 Radians %20
|
||||
%54 = OpExtInst %4 %1 Degrees %22
|
||||
%55 = OpExtInst %4 %1 Radians %22
|
||||
%56 = OpExtInst %4 %1 FClamp %22 %22 %57
|
||||
%58 = OpExtInst %4 %1 Refract %22 %22 %20
|
||||
%59 = OpExtInst %3 %1 Ldexp %20 %39
|
||||
%60 = OpExtInst %10 %1 Ldexp %41 %44
|
||||
%61 = OpExtInst %11 %1 ModfStruct %45
|
||||
%62 = OpExtInst %11 %1 ModfStruct %45
|
||||
%63 = OpCompositeExtract %3 %62 0
|
||||
%64 = OpExtInst %11 %1 ModfStruct %45
|
||||
%65 = OpCompositeExtract %3 %64 1
|
||||
%66 = OpExtInst %12 %1 ModfStruct %46
|
||||
%67 = OpExtInst %13 %1 ModfStruct %47
|
||||
%68 = OpCompositeExtract %4 %67 1
|
||||
%69 = OpCompositeExtract %3 %68 0
|
||||
%70 = OpExtInst %12 %1 ModfStruct %46
|
||||
%71 = OpCompositeExtract %10 %70 0
|
||||
%72 = OpCompositeExtract %3 %71 1
|
||||
%73 = OpExtInst %14 %1 FrexpStruct %45
|
||||
%74 = OpExtInst %14 %1 FrexpStruct %45
|
||||
%75 = OpCompositeExtract %3 %74 0
|
||||
%76 = OpExtInst %14 %1 FrexpStruct %45
|
||||
%77 = OpCompositeExtract %5 %76 1
|
||||
%78 = OpExtInst %15 %1 FrexpStruct %47
|
||||
%79 = OpCompositeExtract %6 %78 1
|
||||
%80 = OpCompositeExtract %5 %79 0
|
||||
%81 = OpQuantizeToF16 %3 %20
|
||||
%82 = OpQuantizeToF16 %10 %48
|
||||
%83 = OpQuantizeToF16 %16 %49
|
||||
%84 = OpQuantizeToF16 %4 %50
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -12,8 +12,7 @@ fn test_integer_dot_product() -> i32 {
|
||||
let a_3_ = vec3(1u);
|
||||
let b_3_ = vec3(1u);
|
||||
let c_3_ = dot(a_3_, b_3_);
|
||||
let c_4_ = dot(vec4(4i), vec4(2i));
|
||||
return c_4_;
|
||||
return 32i;
|
||||
}
|
||||
|
||||
fn test_packed_integer_dot_product() -> u32 {
|
||||
|
||||
@ -9,7 +9,6 @@ fn main() {
|
||||
let g = refract(v, v, 1f);
|
||||
let sign_b = vec4<i32>(-1i, -1i, -1i, -1i);
|
||||
let sign_d = vec4<f32>(-1f, -1f, -1f, -1f);
|
||||
let const_dot = dot(vec2<i32>(), vec2<i32>());
|
||||
let flb_b = vec2<i32>(-1i, -1i);
|
||||
let flb_c = vec2<u32>(0u, 0u);
|
||||
let ftb_c = vec2<i32>(0i, 0i);
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user