From 8fd4ef26f5c6ede28d09d79d52dbdde931d1a718 Mon Sep 17 00:00:00 2001 From: dan sinclair Date: Wed, 5 Oct 2022 17:49:00 +0000 Subject: [PATCH] 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()` it does `int3()`. Bug: tint:1677 Change-Id: I72218c06853e4e5ae1a0d34e2fc3e1ca597de993 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/104682 Kokoro: Kokoro Reviewed-by: Ben Clayton Auto-Submit: Dan Sinclair Commit-Queue: Dan Sinclair Commit-Queue: Ben Clayton --- src/tint/writer/msl/generator_impl.cc | 30 +++++++++++++++++-- test/tint/bug/tint/1121.wgsl.expected.msl | 2 +- test/tint/bug/tint/1520.spvasm.expected.msl | 12 ++++---- test/tint/bug/tint/1677.wgsl.expected.msl | 2 +- test/tint/bug/tint/942.wgsl.expected.msl | 6 ++-- .../add/scalar-vec3/i32.wgsl.expected.msl | 2 +- .../add/vec3-scalar/i32.wgsl.expected.msl | 2 +- .../add/vec3-vec3/i32.wgsl.expected.msl | 2 +- .../scalar-vec3/i32.wgsl.expected.msl | 2 +- .../vec3-vec3/i32.wgsl.expected.msl | 2 +- .../scalar-vec3/i32.wgsl.expected.msl | 2 +- .../vec3-vec3/i32.wgsl.expected.msl | 2 +- .../mul/scalar-vec3/i32.wgsl.expected.msl | 2 +- .../mul/vec3-scalar/i32.wgsl.expected.msl | 2 +- .../mul/vec3-vec3/i32.wgsl.expected.msl | 2 +- .../sub/scalar-vec3/i32.wgsl.expected.msl | 2 +- .../sub/vec3-scalar/i32.wgsl.expected.msl | 2 +- .../sub/vec3-vec3/i32.wgsl.expected.msl | 2 +- .../vector/minus.wgsl.expected.msl | 2 +- .../vector/plus.wgsl.expected.msl | 2 +- .../vector/times.wgsl.expected.msl | 2 +- 21 files changed, 55 insertions(+), 29 deletions(-) diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc index f774315e3d..95b32c0d18 100644 --- a/src/tint/writer/msl/generator_impl.cc +++ b/src/tint/writer/msl/generator_impl.cc @@ -152,6 +152,32 @@ class ScopedBitCast { 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(); + + // 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 @@ bool GeneratorImpl::EmitBinary(std::ostream& out, const ast::BinaryExpression* e // 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 @@ bool GeneratorImpl::EmitBinary(std::ostream& out, const ast::BinaryExpression* e // 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 33fb2df1f4..34c70b2d60 100644 --- a/test/tint/bug/tint/1121.wgsl.expected.msl +++ b/test/tint/bug/tint/1121.wgsl.expected.msl @@ -88,7 +88,7 @@ void tint_symbol_inner(uint3 GlobalInvocationID, const constant Config* const ti for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = as_type((as_type(x_1) + as_type(1)))) { int2 tilePixel0Idx = int2(as_type((as_type(x_1) * as_type(TILE_SIZE))), as_type((as_type(y_1) * as_type(TILE_SIZE)))); float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f)); - float2 ceilCoord = (((2.0f * float2(as_type((as_type(as_type(tilePixel0Idx)) + as_type(as_type(int2(TILE_SIZE))))))) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f)); + float2 ceilCoord = (((2.0f * float2(as_type((as_type(int2(tilePixel0Idx)) + as_type(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 a409d7dcbd..3b8e4f7a59 100644 --- a/test/tint/bug/tint/1520.spvasm.expected.msl +++ b/test/tint/bug/tint/1520.spvasm.expected.msl @@ -45,13 +45,13 @@ bool test_int_S1_c0_b(const constant UniformBuffer* const tint_symbol_5) { 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((as_type(as_type(x_44)) + as_type(as_type(int4(1))))); + int4 const x_47 = as_type((as_type(int4(x_44)) + as_type(int4(int4(1))))); val = x_47; - int4 const x_48 = as_type((as_type(as_type(x_47)) - as_type(as_type(int4(1))))); + int4 const x_48 = as_type((as_type(int4(x_47)) - as_type(int4(int4(1))))); val = x_48; - int4 const x_49 = as_type((as_type(as_type(x_48)) + as_type(as_type(int4(1))))); + int4 const x_49 = as_type((as_type(int4(x_48)) + as_type(int4(int4(1))))); val = x_49; - int4 const x_50 = as_type((as_type(as_type(x_49)) - as_type(as_type(int4(1))))); + int4 const x_50 = as_type((as_type(int4(x_49)) - as_type(int4(int4(1))))); val = x_50; x_55 = false; if (x_41) { @@ -59,11 +59,11 @@ bool test_int_S1_c0_b(const constant UniformBuffer* const tint_symbol_5) { x_55 = x_54; } ok = x_55; - int4 const x_58 = as_type((as_type(as_type(x_50)) * as_type(as_type(int4(2))))); + int4 const x_58 = as_type((as_type(int4(x_50)) * as_type(int4(int4(2))))); val = x_58; int4 const x_59 = (x_58 / int4(2)); val = x_59; - int4 const x_60 = as_type((as_type(as_type(x_59)) * as_type(as_type(int4(2))))); + int4 const x_60 = as_type((as_type(int4(x_59)) * as_type(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 8c3c01957c..6ee8bbda91 100644 --- a/test/tint/bug/tint/1677.wgsl.expected.msl +++ b/test/tint/bug/tint/1677.wgsl.expected.msl @@ -20,7 +20,7 @@ struct Input { }; void tint_symbol_inner(uint3 id, const device Input* const tint_symbol_1) { - int3 const pos = as_type((as_type(as_type((*(tint_symbol_1)).position)) - as_type(as_type(int3(0))))); + int3 const pos = as_type((as_type(int3((*(tint_symbol_1)).position)) - as_type(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 868cd1e697..d605a4f42e 100644 --- a/test/tint/bug/tint/942.wgsl.expected.msl +++ b/test/tint/bug/tint/942.wgsl.expected.msl @@ -32,10 +32,10 @@ void tint_symbol_inner(uint3 WorkGroupID, uint3 LocalInvocationID, uint local_in 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((as_type(as_type(int2(((uint3(WorkGroupID).xy * uint2((*(tint_symbol_2)).blockDim, 4u)) + (uint3(LocalInvocationID).xy * uint2(4u, 1u)))))) - as_type(as_type(int2(int(filterOffset), 0))))); + int2 const baseIndex = as_type((as_type(int2(int2(((uint3(WorkGroupID).xy * uint2((*(tint_symbol_2)).blockDim, 4u)) + (uint3(LocalInvocationID).xy * uint2(4u, 1u)))))) - as_type(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((as_type(as_type(baseIndex)) + as_type(as_type(int2(int(c), int(r)))))); + int2 loadIndex = as_type((as_type(int2(baseIndex)) + as_type(int2(int2(int(c), int(r)))))); if (((*(tint_symbol_4)).value != 0u)) { loadIndex = int2(loadIndex).yx; } @@ -45,7 +45,7 @@ void tint_symbol_inner(uint3 WorkGroupID, uint3 LocalInvocationID, uint local_in 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((as_type(as_type(baseIndex)) + as_type(as_type(int2(int(c), int(r)))))); + int2 writeIndex = as_type((as_type(int2(baseIndex)) + as_type(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 8ba40b7f67..8658c03602 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 @@ using namespace metal; kernel void f() { int const a = 4; int3 const b = int3(1, 2, 3); - int3 const r = as_type((as_type(a) + as_type(as_type(b)))); + int3 const r = as_type((as_type(a) + as_type(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 d5f6b0ec72..de4195f275 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 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); int const b = 4; - int3 const r = as_type((as_type(as_type(a)) + as_type(b))); + int3 const r = as_type((as_type(int3(a)) + as_type(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 a60e42ac54..009e87ebc9 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 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); int3 const b = int3(4, 5, 6); - int3 const r = as_type((as_type(as_type(a)) + as_type(as_type(b)))); + int3 const r = as_type((as_type(int3(a)) + as_type(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 6f38d4bbbd..6d2f8bc123 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 @@ using namespace metal; kernel void f() { int a = 4; int3 b = int3(0, 2, 0); - int3 const r = (a / as_type((as_type(as_type(b)) + as_type(as_type(b))))); + int3 const r = (a / as_type((as_type(int3(b)) + as_type(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 68a2b336c5..c90c00e080 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 @@ using namespace metal; kernel void f() { int3 a = int3(1, 2, 3); int3 b = int3(0, 5, 0); - int3 const r = (a / as_type((as_type(as_type(b)) + as_type(as_type(b))))); + int3 const r = (a / as_type((as_type(int3(b)) + as_type(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 2545d86458..22f6994df9 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 @@ using namespace metal; kernel void f() { int a = 4; int3 b = int3(0, 2, 0); - int3 const r = (a % as_type((as_type(as_type(b)) + as_type(as_type(b))))); + int3 const r = (a % as_type((as_type(int3(b)) + as_type(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 4ed1ec264b..2d83cd38eb 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 @@ using namespace metal; kernel void f() { int3 a = int3(1, 2, 3); int3 b = int3(0, 5, 0); - int3 const r = (a % as_type((as_type(as_type(b)) + as_type(as_type(b))))); + int3 const r = (a % as_type((as_type(int3(b)) + as_type(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 71e20e6ad3..0a492f4854 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 @@ using namespace metal; kernel void f() { int const a = 4; int3 const b = int3(1, 2, 3); - int3 const r = as_type((as_type(a) * as_type(as_type(b)))); + int3 const r = as_type((as_type(a) * as_type(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 8ddccd9a00..134ba04d8f 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 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); int const b = 4; - int3 const r = as_type((as_type(as_type(a)) * as_type(b))); + int3 const r = as_type((as_type(int3(a)) * as_type(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 554bc145e9..dbbcac50ca 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 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); int3 const b = int3(4, 5, 6); - int3 const r = as_type((as_type(as_type(a)) * as_type(as_type(b)))); + int3 const r = as_type((as_type(int3(a)) * as_type(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 2021a7a0cd..5403240cbc 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 @@ using namespace metal; kernel void f() { int const a = 4; int3 const b = int3(1, 2, 3); - int3 const r = as_type((as_type(a) - as_type(as_type(b)))); + int3 const r = as_type((as_type(a) - as_type(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 efe23d7258..9e59fcd91b 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 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); int const b = 4; - int3 const r = as_type((as_type(as_type(a)) - as_type(b))); + int3 const r = as_type((as_type(int3(a)) - as_type(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 2cbef7a881..a42ec6ce13 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 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); int3 const b = int3(4, 5, 6); - int3 const r = as_type((as_type(as_type(a)) - as_type(as_type(b)))); + int3 const r = as_type((as_type(int3(a)) - as_type(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 7984259605..d6ddc129df 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 @@ struct S { }; void foo(device S* const tint_symbol) { - (*(tint_symbol)).a = as_type((as_type(as_type((*(tint_symbol)).a)) - as_type(as_type(int4(2))))); + (*(tint_symbol)).a = as_type((as_type(int4((*(tint_symbol)).a)) - as_type(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 97958782de..455e61d7ea 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 @@ struct S { }; void foo(device S* const tint_symbol) { - (*(tint_symbol)).a = as_type((as_type(as_type((*(tint_symbol)).a)) + as_type(as_type(int4(2))))); + (*(tint_symbol)).a = as_type((as_type(int4((*(tint_symbol)).a)) + as_type(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 9fec5f3ab4..3dfa83f58c 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 @@ struct S { }; void foo(device S* const tint_symbol) { - (*(tint_symbol)).a = as_type((as_type(as_type((*(tint_symbol)).a)) * as_type(as_type(int4(2))))); + (*(tint_symbol)).a = as_type((as_type(int4((*(tint_symbol)).a)) * as_type(int4(int4(2))))); }