Fix MSL packed_int casting.
In 104681 the vectors were cast to themselves to fixup an issue
with `packed_int`. That CL used an `as_type` which does a bit cast.
A `packed_int` can not be bitcast to an `int`. This CL changes to
a type cast, so instead of `as_type<int3>()` it does `int3()`.
Bug: tint:1677
Change-Id: I72218c06853e4e5ae1a0d34e2fc3e1ca597de993
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/104682
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Auto-Submit: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index f774315..95b32c0 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -152,6 +152,32 @@
std::ostream& s;
};
+class ScopedCast {
+ public:
+ ScopedCast(GeneratorImpl* generator,
+ std::ostream& stream,
+ const sem::Type* curr_type,
+ const sem::Type* target_type)
+ : s(stream) {
+ auto* target_vec_type = target_type->As<sem::Vector>();
+
+ // If we need to promote from scalar to vector, cast the scalar to the
+ // vector element type.
+ if (curr_type->is_scalar() && target_vec_type) {
+ target_type = target_vec_type->type();
+ }
+
+ // Cast
+ generator->EmitType(s, target_type, "");
+ s << "(";
+ }
+
+ ~ScopedCast() { s << ")"; }
+
+ private:
+ std::ostream& s;
+};
+
} // namespace
SanitizedResult::SanitizedResult() = default;
@@ -521,7 +547,7 @@
// In case the type is packed, cast to our own type in order to remove the packing.
// Otherwise, this just casts to itself.
if (lhs_type->is_signed_integer_vector()) {
- ScopedBitCast lhs_self_cast(this, out, lhs_type, lhs_type);
+ ScopedCast lhs_self_cast(this, out, lhs_type, lhs_type);
if (!EmitExpression(out, expr->lhs)) {
return false;
}
@@ -540,7 +566,7 @@
// In case the type is packed, cast to our own type in order to remove the packing.
// Otherwise, this just casts to itself.
if (rhs_type->is_signed_integer_vector()) {
- ScopedBitCast rhs_self_cast(this, out, rhs_type, rhs_type);
+ ScopedCast rhs_self_cast(this, out, rhs_type, rhs_type);
if (!EmitExpression(out, expr->rhs)) {
return false;
}
diff --git a/test/tint/bug/tint/1121.wgsl.expected.msl b/test/tint/bug/tint/1121.wgsl.expected.msl
index 33fb2df..34c70b2 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.msl
+++ b/test/tint/bug/tint/1121.wgsl.expected.msl
@@ -88,7 +88,7 @@
for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = as_type<int>((as_type<uint>(x_1) + as_type<uint>(1)))) {
int2 tilePixel0Idx = int2(as_type<int>((as_type<uint>(x_1) * as_type<uint>(TILE_SIZE))), as_type<int>((as_type<uint>(y_1) * as_type<uint>(TILE_SIZE))));
float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
- float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(as_type<int2>(tilePixel0Idx)) + as_type<uint2>(as_type<int2>(int2(TILE_SIZE))))))) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
+ float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(int2(tilePixel0Idx)) + as_type<uint2>(int2(int2(TILE_SIZE))))))) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
float2 viewFloorCoord = float2((((-(viewNear) * floorCoord[0]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord[1]) - (M[2][1] * viewNear)) / M[1][1]));
float2 viewCeilCoord = float2((((-(viewNear) * ceilCoord[0]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord[1]) - (M[2][1] * viewNear)) / M[1][1]));
frustumPlanes[0] = float4(1.0f, 0.0f, (-(viewFloorCoord[0]) / viewNear), 0.0f);
diff --git a/test/tint/bug/tint/1520.spvasm.expected.msl b/test/tint/bug/tint/1520.spvasm.expected.msl
index a409d7d..3b8e4f7 100644
--- a/test/tint/bug/tint/1520.spvasm.expected.msl
+++ b/test/tint/bug/tint/1520.spvasm.expected.msl
@@ -45,13 +45,13 @@
ok = x_41;
int4 const x_44 = int4(x_27, x_27, x_27, x_27);
val = x_44;
- int4 const x_47 = as_type<int4>((as_type<uint4>(as_type<int4>(x_44)) + as_type<uint4>(as_type<int4>(int4(1)))));
+ int4 const x_47 = as_type<int4>((as_type<uint4>(int4(x_44)) + as_type<uint4>(int4(int4(1)))));
val = x_47;
- int4 const x_48 = as_type<int4>((as_type<uint4>(as_type<int4>(x_47)) - as_type<uint4>(as_type<int4>(int4(1)))));
+ int4 const x_48 = as_type<int4>((as_type<uint4>(int4(x_47)) - as_type<uint4>(int4(int4(1)))));
val = x_48;
- int4 const x_49 = as_type<int4>((as_type<uint4>(as_type<int4>(x_48)) + as_type<uint4>(as_type<int4>(int4(1)))));
+ int4 const x_49 = as_type<int4>((as_type<uint4>(int4(x_48)) + as_type<uint4>(int4(int4(1)))));
val = x_49;
- int4 const x_50 = as_type<int4>((as_type<uint4>(as_type<int4>(x_49)) - as_type<uint4>(as_type<int4>(int4(1)))));
+ int4 const x_50 = as_type<int4>((as_type<uint4>(int4(x_49)) - as_type<uint4>(int4(int4(1)))));
val = x_50;
x_55 = false;
if (x_41) {
@@ -59,11 +59,11 @@
x_55 = x_54;
}
ok = x_55;
- int4 const x_58 = as_type<int4>((as_type<uint4>(as_type<int4>(x_50)) * as_type<uint4>(as_type<int4>(int4(2)))));
+ int4 const x_58 = as_type<int4>((as_type<uint4>(int4(x_50)) * as_type<uint4>(int4(int4(2)))));
val = x_58;
int4 const x_59 = (x_58 / int4(2));
val = x_59;
- int4 const x_60 = as_type<int4>((as_type<uint4>(as_type<int4>(x_59)) * as_type<uint4>(as_type<int4>(int4(2)))));
+ int4 const x_60 = as_type<int4>((as_type<uint4>(int4(x_59)) * as_type<uint4>(int4(int4(2)))));
val = x_60;
int4 const x_61 = (x_60 / int4(2));
val = x_61;
diff --git a/test/tint/bug/tint/1677.wgsl.expected.msl b/test/tint/bug/tint/1677.wgsl.expected.msl
index 8c3c019..6ee8bbd 100644
--- a/test/tint/bug/tint/1677.wgsl.expected.msl
+++ b/test/tint/bug/tint/1677.wgsl.expected.msl
@@ -20,7 +20,7 @@
};
void tint_symbol_inner(uint3 id, const device Input* const tint_symbol_1) {
- int3 const pos = as_type<int3>((as_type<uint3>(as_type<int3>((*(tint_symbol_1)).position)) - as_type<uint3>(as_type<int3>(int3(0)))));
+ int3 const pos = as_type<int3>((as_type<uint3>(int3((*(tint_symbol_1)).position)) - as_type<uint3>(int3(int3(0)))));
}
kernel void tint_symbol(const device Input* tint_symbol_2 [[buffer(0)]], uint3 id [[thread_position_in_grid]]) {
diff --git a/test/tint/bug/tint/942.wgsl.expected.msl b/test/tint/bug/tint/942.wgsl.expected.msl
index 868cd1e..d605a4f 100644
--- a/test/tint/bug/tint/942.wgsl.expected.msl
+++ b/test/tint/bug/tint/942.wgsl.expected.msl
@@ -32,10 +32,10 @@
threadgroup_barrier(mem_flags::mem_threadgroup);
uint const filterOffset = (((*(tint_symbol_2)).filterDim - 1u) / 2u);
int2 const dims = int2(tint_symbol_3.get_width(0), tint_symbol_3.get_height(0));
- int2 const baseIndex = as_type<int2>((as_type<uint2>(as_type<int2>(int2(((uint3(WorkGroupID).xy * uint2((*(tint_symbol_2)).blockDim, 4u)) + (uint3(LocalInvocationID).xy * uint2(4u, 1u)))))) - as_type<uint2>(as_type<int2>(int2(int(filterOffset), 0)))));
+ int2 const baseIndex = as_type<int2>((as_type<uint2>(int2(int2(((uint3(WorkGroupID).xy * uint2((*(tint_symbol_2)).blockDim, 4u)) + (uint3(LocalInvocationID).xy * uint2(4u, 1u)))))) - as_type<uint2>(int2(int2(int(filterOffset), 0)))));
for(uint r = 0u; (r < 4u); r = (r + 1u)) {
for(uint c = 0u; (c < 4u); c = (c + 1u)) {
- int2 loadIndex = as_type<int2>((as_type<uint2>(as_type<int2>(baseIndex)) + as_type<uint2>(as_type<int2>(int2(int(c), int(r))))));
+ int2 loadIndex = as_type<int2>((as_type<uint2>(int2(baseIndex)) + as_type<uint2>(int2(int2(int(c), int(r))))));
if (((*(tint_symbol_4)).value != 0u)) {
loadIndex = int2(loadIndex).yx;
}
@@ -45,7 +45,7 @@
threadgroup_barrier(mem_flags::mem_threadgroup);
for(uint r = 0u; (r < 4u); r = (r + 1u)) {
for(uint c = 0u; (c < 4u); c = (c + 1u)) {
- int2 writeIndex = as_type<int2>((as_type<uint2>(as_type<int2>(baseIndex)) + as_type<uint2>(as_type<int2>(int2(int(c), int(r))))));
+ int2 writeIndex = as_type<int2>((as_type<uint2>(int2(baseIndex)) + as_type<uint2>(int2(int2(int(c), int(r))))));
if (((*(tint_symbol_4)).value != 0u)) {
writeIndex = int2(writeIndex).yx;
}
diff --git a/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl
index 8ba40b7..8658c03 100644
--- a/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
kernel void f() {
int const a = 4;
int3 const b = int3(1, 2, 3);
- int3 const r = as_type<int3>((as_type<uint>(a) + as_type<uint3>(as_type<int3>(b))));
+ int3 const r = as_type<int3>((as_type<uint>(a) + as_type<uint3>(int3(b))));
return;
}
diff --git a/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl b/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl
index d5f6b0e..de4195f 100644
--- a/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
kernel void f() {
int3 const a = int3(1, 2, 3);
int const b = 4;
- int3 const r = as_type<int3>((as_type<uint3>(as_type<int3>(a)) + as_type<uint>(b)));
+ int3 const r = as_type<int3>((as_type<uint3>(int3(a)) + as_type<uint>(b)));
return;
}
diff --git a/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl
index a60e42a..009e87e 100644
--- a/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
kernel void f() {
int3 const a = int3(1, 2, 3);
int3 const b = int3(4, 5, 6);
- int3 const r = as_type<int3>((as_type<uint3>(as_type<int3>(a)) + as_type<uint3>(as_type<int3>(b))));
+ int3 const r = as_type<int3>((as_type<uint3>(int3(a)) + as_type<uint3>(int3(b))));
return;
}
diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
index 6f38d4b..6d2f8bc 100644
--- a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
kernel void f() {
int a = 4;
int3 b = int3(0, 2, 0);
- int3 const r = (a / as_type<int3>((as_type<uint3>(as_type<int3>(b)) + as_type<uint3>(as_type<int3>(b)))));
+ int3 const r = (a / as_type<int3>((as_type<uint3>(int3(b)) + as_type<uint3>(int3(b)))));
return;
}
diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
index 68a2b33..c90c00e 100644
--- a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
kernel void f() {
int3 a = int3(1, 2, 3);
int3 b = int3(0, 5, 0);
- int3 const r = (a / as_type<int3>((as_type<uint3>(as_type<int3>(b)) + as_type<uint3>(as_type<int3>(b)))));
+ int3 const r = (a / as_type<int3>((as_type<uint3>(int3(b)) + as_type<uint3>(int3(b)))));
return;
}
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
index 2545d86..22f6994 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
kernel void f() {
int a = 4;
int3 b = int3(0, 2, 0);
- int3 const r = (a % as_type<int3>((as_type<uint3>(as_type<int3>(b)) + as_type<uint3>(as_type<int3>(b)))));
+ int3 const r = (a % as_type<int3>((as_type<uint3>(int3(b)) + as_type<uint3>(int3(b)))));
return;
}
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
index 4ed1ec2..2d83cd3 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
kernel void f() {
int3 a = int3(1, 2, 3);
int3 b = int3(0, 5, 0);
- int3 const r = (a % as_type<int3>((as_type<uint3>(as_type<int3>(b)) + as_type<uint3>(as_type<int3>(b)))));
+ int3 const r = (a % as_type<int3>((as_type<uint3>(int3(b)) + as_type<uint3>(int3(b)))));
return;
}
diff --git a/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl
index 71e20e6..0a492f4 100644
--- a/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
kernel void f() {
int const a = 4;
int3 const b = int3(1, 2, 3);
- int3 const r = as_type<int3>((as_type<uint>(a) * as_type<uint3>(as_type<int3>(b))));
+ int3 const r = as_type<int3>((as_type<uint>(a) * as_type<uint3>(int3(b))));
return;
}
diff --git a/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl b/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl
index 8ddccd9..134ba04 100644
--- a/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
kernel void f() {
int3 const a = int3(1, 2, 3);
int const b = 4;
- int3 const r = as_type<int3>((as_type<uint3>(as_type<int3>(a)) * as_type<uint>(b)));
+ int3 const r = as_type<int3>((as_type<uint3>(int3(a)) * as_type<uint>(b)));
return;
}
diff --git a/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl
index 554bc14..dbbcac5 100644
--- a/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
kernel void f() {
int3 const a = int3(1, 2, 3);
int3 const b = int3(4, 5, 6);
- int3 const r = as_type<int3>((as_type<uint3>(as_type<int3>(a)) * as_type<uint3>(as_type<int3>(b))));
+ int3 const r = as_type<int3>((as_type<uint3>(int3(a)) * as_type<uint3>(int3(b))));
return;
}
diff --git a/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl
index 2021a7a..5403240 100644
--- a/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
kernel void f() {
int const a = 4;
int3 const b = int3(1, 2, 3);
- int3 const r = as_type<int3>((as_type<uint>(a) - as_type<uint3>(as_type<int3>(b))));
+ int3 const r = as_type<int3>((as_type<uint>(a) - as_type<uint3>(int3(b))));
return;
}
diff --git a/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl b/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl
index efe23d7..9e59fcd 100644
--- a/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
kernel void f() {
int3 const a = int3(1, 2, 3);
int const b = 4;
- int3 const r = as_type<int3>((as_type<uint3>(as_type<int3>(a)) - as_type<uint>(b)));
+ int3 const r = as_type<int3>((as_type<uint3>(int3(a)) - as_type<uint>(b)));
return;
}
diff --git a/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl
index 2cbef7a..a42ec6c 100644
--- a/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
kernel void f() {
int3 const a = int3(1, 2, 3);
int3 const b = int3(4, 5, 6);
- int3 const r = as_type<int3>((as_type<uint3>(as_type<int3>(a)) - as_type<uint3>(as_type<int3>(b))));
+ int3 const r = as_type<int3>((as_type<uint3>(int3(a)) - as_type<uint3>(int3(b))));
return;
}
diff --git a/test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl
index 7984259..d6ddc12 100644
--- a/test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl
@@ -6,6 +6,6 @@
};
void foo(device S* const tint_symbol) {
- (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(as_type<int4>((*(tint_symbol)).a)) - as_type<uint4>(as_type<int4>(int4(2)))));
+ (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(int4((*(tint_symbol)).a)) - as_type<uint4>(int4(int4(2)))));
}
diff --git a/test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl
index 9795878..455e61d 100644
--- a/test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl
@@ -6,6 +6,6 @@
};
void foo(device S* const tint_symbol) {
- (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(as_type<int4>((*(tint_symbol)).a)) + as_type<uint4>(as_type<int4>(int4(2)))));
+ (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(int4((*(tint_symbol)).a)) + as_type<uint4>(int4(int4(2)))));
}
diff --git a/test/tint/statements/compound_assign/vector/times.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/times.wgsl.expected.msl
index 9fec5f3..3dfa83f 100644
--- a/test/tint/statements/compound_assign/vector/times.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/vector/times.wgsl.expected.msl
@@ -6,6 +6,6 @@
};
void foo(device S* const tint_symbol) {
- (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(as_type<int4>((*(tint_symbol)).a)) * as_type<uint4>(as_type<int4>(int4(2)))));
+ (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(int4((*(tint_symbol)).a)) * as_type<uint4>(int4(int4(2)))));
}
