github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

wgsl-in: Splat on compound assignments

Browse Source 
Compound assignments on wgsl follow the same semantics as their
underlying operation, this includes the splatting behavior when mixing
scalar and vector operands, which was done for binary operations but not
for compound assignments.
This commit is contained in:
João Capucho
2022-09-07 22:58:40 +01:00
committed by Jim Blandy
parent eb6124ec94
commit db1321cfb6
7 changed files with 711 additions and 605 deletions
Download Patch File Download Diff File Expand all files Collapse all files

67
src/front/wgsl/mod.rs
View File

@@ -916,29 +916,8 @@ impl<'a> ExpressionContext<'a, '_, '_> {
let mut right = self.apply_load_rule(unloaded_right);
let end = lexer.current_byte_offset() as u32;
// Insert splats, if needed by the non-'*' operations.
// (`BinaryOperator::Multiply` handles splats itself.)
if op != crate::BinaryOperator::Multiply {
let left_size = match *self.resolve_type(left)? {
crate::TypeInner::Vector { size, .. } => Some(size),
_ => None,
};
match (left_size, self.resolve_type(right)?) {
(Some(size), &crate::TypeInner::Scalar { .. }) => {
right = self.expressions.append(
crate::Expression::Splat { size, value: right },
self.expressions.get_span(right),
);
}
(None, &crate::TypeInner::Vector { size, .. }) => {
left = self.expressions.append(
crate::Expression::Splat { size, value: left },
self.expressions.get_span(left),
);
}
_ => {}
}
}
self.binary_op_splat(op, &mut left, &mut right)?;
accumulator = TypedExpression::non_reference(self.expressions.append(
crate::Expression::Binary { op, left, right },
NagaSpan::new(start, end),
@@ -947,6 +926,41 @@ impl<'a> ExpressionContext<'a, '_, '_> {
Ok(accumulator)
}
/// Insert splats, if needed by the non-'*' operations.
fn binary_op_splat(
&mut self,
op: crate::BinaryOperator,
left: &mut Handle<crate::Expression>,
right: &mut Handle<crate::Expression>,
) -> Result<(), Error<'a>> {
if op != crate::BinaryOperator::Multiply {
let left_size = match *self.resolve_type(*left)? {
crate::TypeInner::Vector { size, .. } => Some(size),
_ => None,
};
match (left_size, self.resolve_type(*right)?) {
(Some(size), &crate::TypeInner::Scalar { .. }) => {
*right = self.expressions.append(
crate::Expression::Splat {
size,
value: *right,
},
self.expressions.get_span(*right),
);
}
(None, &crate::TypeInner::Vector { size, .. }) => {
*left = self.expressions.append(
crate::Expression::Splat { size, value: *left },
self.expressions.get_span(*left),
);
}
_ => {}
}
}
Ok(())
}
/// Add a single expression to the expression table that is not covered by `self.emitter`.
///
/// This is useful for `CallResult` and `AtomicResult` expressions, which should not be covered by
@@ -3321,13 +3335,16 @@ impl Parser {
//Note: `consume_token` shouldn't produce any other assignment ops
_ => unreachable!(),
};
let left = context.expressions.append(
let mut left = context.expressions.append(
crate::Expression::Load {
pointer: reference.handle,
},
lhs_span.into(),
);
let right = self.parse_general_expression(lexer, context.reborrow())?;
let mut right = self.parse_general_expression(lexer, context.reborrow())?;
context.binary_op_splat(op, &mut left, &mut right)?;
context
.expressions
.append(crate::Expression::Binary { op, left, right }, span.into())

8
tests/in/operators.wgsl
View File

@@ -28,6 +28,14 @@ fn splat() -> vec4<f32> {
return a.xyxy + vec4<f32>(b);
}
fn splat_assignment() -> vec2<f32> {
var a = vec2<f32>(2.0);
a += 1.0;
a -= 3.0;
a /= 4.0;
return a;
}
fn bool_cast(x: vec3<f32>) -> vec3<f32> {
let y = vec3<bool>(x);
return vec3<f32>(y);

71
tests/out/glsl/operators.main.Compute.glsl
View File

@@ -23,9 +23,22 @@ vec4 builtins() {
}
vec4 splat() {
vec2 a_1 = (((vec2(1.0) + vec2(2.0)) - vec2(3.0)) / vec2(4.0));
vec2 a_2 = (((vec2(1.0) + vec2(2.0)) - vec2(3.0)) / vec2(4.0));
ivec4 b = (ivec4(5) % ivec4(2));
return (a_1.xyxy + vec4(b));
return (a_2.xyxy + vec4(b));
}
vec2 splat_assignment() {
vec2 a = vec2(0.0);
a = vec2(2.0);
vec2 _e7 = a;
a = (_e7 + vec2(1.0));
vec2 _e11 = a;
a = (_e11 - vec2(3.0));
vec2 _e15 = a;
a = (_e15 / vec2(4.0));
vec2 _e19 = a;
return _e19;
}
vec3 bool_cast(vec3 x) {
@@ -203,34 +216,34 @@ void comparison() {
}
void assignment() {
int a = 1;
int a_1 = 1;
ivec3 vec0_ = ivec3(0, 0, 0);
int _e6 = a;
a = (_e6 + 1);
int _e9 = a;
a = (_e9 - 1);
int _e12 = a;
int _e13 = a;
a = (_e12 * _e13);
int _e15 = a;
int _e16 = a;
a = (_e15 / _e16);
int _e18 = a;
a = (_e18 % 1);
int _e21 = a;
a = (_e21 & 0);
int _e24 = a;
a = (_e24 | 0);
int _e27 = a;
a = (_e27 ^ 0);
int _e30 = a;
a = (_e30 << 2u);
int _e33 = a;
a = (_e33 >> 1u);
int _e36 = a;
a = (_e36 + 1);
int _e39 = a;
a = (_e39 - 1);
int _e6 = a_1;
a_1 = (_e6 + 1);
int _e9 = a_1;
a_1 = (_e9 - 1);
int _e12 = a_1;
int _e13 = a_1;
a_1 = (_e12 * _e13);
int _e15 = a_1;
int _e16 = a_1;
a_1 = (_e15 / _e16);
int _e18 = a_1;
a_1 = (_e18 % 1);
int _e21 = a_1;
a_1 = (_e21 & 0);
int _e24 = a_1;
a_1 = (_e24 | 0);
int _e27 = a_1;
a_1 = (_e27 ^ 0);
int _e30 = a_1;
a_1 = (_e30 << 2u);
int _e33 = a_1;
a_1 = (_e33 >> 1u);
int _e36 = a_1;
a_1 = (_e36 + 1);
int _e39 = a_1;
a_1 = (_e39 - 1);
int _e46 = vec0_.y;
vec0_.y = (_e46 + 1);
int _e51 = vec0_.y;

73
tests/out/hlsl/operators.hlsl
View File

@@ -39,9 +39,24 @@ float4 builtins()
float4 splat()
{
float2 a_1 = ((((1.0).xx + (2.0).xx) - (3.0).xx) / (4.0).xx);
float2 a_2 = ((((1.0).xx + (2.0).xx) - (3.0).xx) / (4.0).xx);
int4 b = ((5).xxxx % (2).xxxx);
return (a_1.xyxy + float4(b));
return (a_2.xyxy + float4(b));
}
float2 splat_assignment()
{
float2 a = (float2)0;
a = (2.0).xx;
float2 _expr7 = a;
a = (_expr7 + (1.0).xx);
float2 _expr11 = a;
a = (_expr11 - (3.0).xx);
float2 _expr15 = a;
a = (_expr15 / (4.0).xx);
float2 _expr19 = a;
return _expr19;
}
float3 bool_cast(float3 x)
@@ -233,35 +248,35 @@ void comparison()
void assignment()
{
int a = 1;
int a_1 = 1;
int3 vec0_ = int3(0, 0, 0);
int _expr6 = a;
a = (_expr6 + 1);
int _expr9 = a;
a = (_expr9 - 1);
int _expr12 = a;
int _expr13 = a;
a = (_expr12 * _expr13);
int _expr15 = a;
int _expr16 = a;
a = (_expr15 / _expr16);
int _expr18 = a;
a = (_expr18 % 1);
int _expr21 = a;
a = (_expr21 & 0);
int _expr24 = a;
a = (_expr24 | 0);
int _expr27 = a;
a = (_expr27 ^ 0);
int _expr30 = a;
a = (_expr30 << 2u);
int _expr33 = a;
a = (_expr33 >> 1u);
int _expr36 = a;
a = (_expr36 + 1);
int _expr39 = a;
a = (_expr39 - 1);
int _expr6 = a_1;
a_1 = (_expr6 + 1);
int _expr9 = a_1;
a_1 = (_expr9 - 1);
int _expr12 = a_1;
int _expr13 = a_1;
a_1 = (_expr12 * _expr13);
int _expr15 = a_1;
int _expr16 = a_1;
a_1 = (_expr15 / _expr16);
int _expr18 = a_1;
a_1 = (_expr18 % 1);
int _expr21 = a_1;
a_1 = (_expr21 & 0);
int _expr24 = a_1;
a_1 = (_expr24 | 0);
int _expr27 = a_1;
a_1 = (_expr27 ^ 0);
int _expr30 = a_1;
a_1 = (_expr30 << 2u);
int _expr33 = a_1;
a_1 = (_expr33 >> 1u);
int _expr36 = a_1;
a_1 = (_expr36 + 1);
int _expr39 = a_1;
a_1 = (_expr39 - 1);
int _expr46 = vec0_.y;
vec0_.y = (_expr46 + 1);
int _expr51 = vec0_.y;

84
tests/out/msl/operators.msl
View File

@@ -19,15 +19,15 @@ constant metal::float4 v_f32_zero = {0.0, 0.0, 0.0, 0.0};
constant metal::float4 v_f32_half = {0.5, 0.5, 0.5, 0.5};
constant metal::int4 v_i32_one = {1, 1, 1, 1};
constant metal::uint2 const_type_11_ = {0u, 0u};
constant metal::float2 const_type_6_ = {0.0, 0.0};
constant metal::float2x2 const_type_7_ = {const_type_6_, const_type_6_};
constant metal::float2 const_type_4_ = {0.0, 0.0};
constant metal::float2x2 const_type_7_ = {const_type_4_, const_type_4_};
constant metal::float4 const_type = {0.0, 0.0, 0.0, 0.0};
constant Foo const_Foo = {const_type, 0};
constant type_12 const_type_12_ = {const_Foo, const_Foo, const_Foo};
constant metal::float3 const_type_4_ = {0.0, 0.0, 0.0};
constant metal::float2x3 const_type_14_ = {const_type_4_, const_type_4_};
constant metal::float3x3 const_type_15_ = {const_type_4_, const_type_4_, const_type_4_};
constant metal::float4x3 const_type_16_ = {const_type_4_, const_type_4_, const_type_4_, const_type_4_};
constant metal::float3 const_type_5_ = {0.0, 0.0, 0.0};
constant metal::float2x3 const_type_14_ = {const_type_5_, const_type_5_};
constant metal::float3x3 const_type_15_ = {const_type_5_, const_type_5_, const_type_5_};
constant metal::float4x3 const_type_16_ = {const_type_5_, const_type_5_, const_type_5_, const_type_5_};
constant metal::float3x4 const_type_17_ = {const_type, const_type, const_type};
constant metal::int3 const_type_18_ = {0, 0, 0};
@@ -46,9 +46,23 @@ metal::float4 builtins(
metal::float4 splat(
) {
metal::float2 a_1 = ((metal::float2(1.0) + metal::float2(2.0)) - metal::float2(3.0)) / metal::float2(4.0);
metal::float2 a_2 = ((metal::float2(1.0) + metal::float2(2.0)) - metal::float2(3.0)) / metal::float2(4.0);
metal::int4 b = metal::int4(5) % metal::int4(2);
return a_1.xyxy + static_cast<metal::float4>(b);
return a_2.xyxy + static_cast<metal::float4>(b);
}
metal::float2 splat_assignment(
) {
metal::float2 a = {};
a = metal::float2(2.0);
metal::float2 _e7 = a;
a = _e7 + metal::float2(1.0);
metal::float2 _e11 = a;
a = _e11 - metal::float2(3.0);
metal::float2 _e15 = a;
a = _e15 / metal::float2(4.0);
metal::float2 _e19 = a;
return _e19;
}
metal::float3 bool_cast(
@@ -234,34 +248,34 @@ void comparison(
void assignment(
) {
int a = 1;
int a_1 = 1;
metal::int3 vec0_ = const_type_18_;
int _e6 = a;
a = _e6 + 1;
int _e9 = a;
a = _e9 - 1;
int _e12 = a;
int _e13 = a;
a = _e12 * _e13;
int _e15 = a;
int _e16 = a;
a = _e15 / _e16;
int _e18 = a;
a = _e18 % 1;
int _e21 = a;
a = _e21 & 0;
int _e24 = a;
a = _e24 | 0;
int _e27 = a;
a = _e27 ^ 0;
int _e30 = a;
a = _e30 << 2u;
int _e33 = a;
a = _e33 >> 1u;
int _e36 = a;
a = _e36 + 1;
int _e39 = a;
a = _e39 - 1;
int _e6 = a_1;
a_1 = _e6 + 1;
int _e9 = a_1;
a_1 = _e9 - 1;
int _e12 = a_1;
int _e13 = a_1;
a_1 = _e12 * _e13;
int _e15 = a_1;
int _e16 = a_1;
a_1 = _e15 / _e16;
int _e18 = a_1;
a_1 = _e18 % 1;
int _e21 = a_1;
a_1 = _e21 & 0;
int _e24 = a_1;
a_1 = _e24 | 0;
int _e27 = a_1;
a_1 = _e27 ^ 0;
int _e30 = a_1;
a_1 = _e30 << 2u;
int _e33 = a_1;
a_1 = _e33 >> 1u;
int _e36 = a_1;
a_1 = _e36 + 1;
int _e39 = a_1;
a_1 = _e39 - 1;
int _e46 = vec0_.y;
vec0_.y = _e46 + 1;
int _e51 = vec0_.y;

941
tests/out/spv/operators.spvasm
View File

File diff suppressed because it is too large Load Diff

72
tests/out/wgsl/operators.wgsl
View File

@@ -20,9 +20,23 @@ fn builtins() -> vec4<f32> {
}
fn splat() -> vec4<f32> {
let a_1 = (((vec2<f32>(1.0) + vec2<f32>(2.0)) - vec2<f32>(3.0)) / vec2<f32>(4.0));
let a_2 = (((vec2<f32>(1.0) + vec2<f32>(2.0)) - vec2<f32>(3.0)) / vec2<f32>(4.0));
let b = (vec4<i32>(5) % vec4<i32>(2));
return (a_1.xyxy + vec4<f32>(b));
return (a_2.xyxy + vec4<f32>(b));
}
fn splat_assignment() -> vec2<f32> {
var a: vec2<f32>;
a = vec2<f32>(2.0);
let _e7 = a;
a = (_e7 + vec2<f32>(1.0));
let _e11 = a;
a = (_e11 - vec2<f32>(3.0));
let _e15 = a;
a = (_e15 / vec2<f32>(4.0));
let _e19 = a;
return _e19;
}
fn bool_cast(x: vec3<f32>) -> vec3<f32> {
@@ -201,35 +215,35 @@ fn comparison() {
}
fn assignment() {
var a: i32 = 1;
var a_1: i32 = 1;
var vec0_: vec3<i32> = vec3<i32>(0, 0, 0);
let _e6 = a;
a = (_e6 + 1);
let _e9 = a;
a = (_e9 - 1);
let _e12 = a;
let _e13 = a;
a = (_e12 * _e13);
let _e15 = a;
let _e16 = a;
a = (_e15 / _e16);
let _e18 = a;
a = (_e18 % 1);
let _e21 = a;
a = (_e21 & 0);
let _e24 = a;
a = (_e24 | 0);
let _e27 = a;
a = (_e27 ^ 0);
let _e30 = a;
a = (_e30 << 2u);
let _e33 = a;
a = (_e33 >> 1u);
let _e36 = a;
a = (_e36 + 1);
let _e39 = a;
a = (_e39 - 1);
let _e6 = a_1;
a_1 = (_e6 + 1);
let _e9 = a_1;
a_1 = (_e9 - 1);
let _e12 = a_1;
let _e13 = a_1;
a_1 = (_e12 * _e13);
let _e15 = a_1;
let _e16 = a_1;
a_1 = (_e15 / _e16);
let _e18 = a_1;
a_1 = (_e18 % 1);
let _e21 = a_1;
a_1 = (_e21 & 0);
let _e24 = a_1;
a_1 = (_e24 | 0);
let _e27 = a_1;
a_1 = (_e27 ^ 0);
let _e30 = a_1;
a_1 = (_e30 << 2u);
let _e33 = a_1;
a_1 = (_e33 >> 1u);
let _e36 = a_1;
a_1 = (_e36 + 1);
let _e39 = a_1;
a_1 = (_e39 - 1);
let _e46 = vec0_.y;
vec0_.y = (_e46 + 1);
let _e51 = vec0_.y;