encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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>
This commit is contained in:
dan sinclair 2022-10-05 17:49:00 +00:00 committed by Dawn LUCI CQ
parent f35c719054
commit 8fd4ef26f5
21 changed files with 55 additions and 29 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

30
src/tint/writer/msl/generator_impl.cc
View File

@ -152,6 +152,32 @@ class ScopedBitCast {
std::ostream& s; 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 } // namespace
SanitizedResult::SanitizedResult() = default; 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. // In case the type is packed, cast to our own type in order to remove the packing.
// Otherwise, this just casts to itself. // Otherwise, this just casts to itself.
if (lhs_type->is_signed_integer_vector()) { 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)) { if (!EmitExpression(out, expr->lhs)) {
return false; 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. // In case the type is packed, cast to our own type in order to remove the packing.
// Otherwise, this just casts to itself. // Otherwise, this just casts to itself.
if (rhs_type->is_signed_integer_vector()) { 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)) { if (!EmitExpression(out, expr->rhs)) {
return false; return false;
} }

2
test/tint/bug/tint/1121.wgsl.expected.msl
View File

@ -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<int>((as_type<uint>(x_1) + as_type<uint>(1)))) { 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)))); 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 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 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])); 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); frustumPlanes[0] = float4(1.0f, 0.0f, (-(viewFloorCoord[0]) / viewNear), 0.0f);

12
test/tint/bug/tint/1520.spvasm.expected.msl
View File

@ -45,13 +45,13 @@ bool test_int_S1_c0_b(const constant UniformBuffer* const tint_symbol_5) {
ok = x_41; ok = x_41;
int4 const x_44 = int4(x_27, x_27, x_27, x_27); int4 const x_44 = int4(x_27, x_27, x_27, x_27);
val = x_44; 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; 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; 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; 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; val = x_50;
x_55 = false; x_55 = false;
if (x_41) { if (x_41) {
@ -59,11 +59,11 @@ bool test_int_S1_c0_b(const constant UniformBuffer* const tint_symbol_5) {
x_55 = x_54; x_55 = x_54;
} }
ok = x_55; 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; val = x_58;
int4 const x_59 = (x_58 / int4(2)); int4 const x_59 = (x_58 / int4(2));
val = x_59; 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; val = x_60;
int4 const x_61 = (x_60 / int4(2)); int4 const x_61 = (x_60 / int4(2));
val = x_61; val = x_61;

2
test/tint/bug/tint/1677.wgsl.expected.msl
View File

@ -20,7 +20,7 @@ struct Input {
}; };
void tint_symbol_inner(uint3 id, const device Input* const tint_symbol_1) { 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]]) { kernel void tint_symbol(const device Input* tint_symbol_2 [[buffer(0)]], uint3 id [[thread_position_in_grid]]) {

6
test/tint/bug/tint/942.wgsl.expected.msl
View File

@ -32,10 +32,10 @@ void tint_symbol_inner(uint3 WorkGroupID, uint3 LocalInvocationID, uint local_in
threadgroup_barrier(mem_flags::mem_threadgroup); threadgroup_barrier(mem_flags::mem_threadgroup);
uint const filterOffset = (((*(tint_symbol_2)).filterDim - 1u) / 2u); 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 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 r = 0u; (r < 4u); r = (r + 1u)) {
for(uint c = 0u; (c < 4u); c = (c + 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)) { if (((*(tint_symbol_4)).value != 0u)) {
loadIndex = int2(loadIndex).yx; 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); threadgroup_barrier(mem_flags::mem_threadgroup);
for(uint r = 0u; (r < 4u); r = (r + 1u)) { for(uint r = 0u; (r < 4u); r = (r + 1u)) {
for(uint c = 0u; (c < 4u); c = (c + 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)) { if (((*(tint_symbol_4)).value != 0u)) {
writeIndex = int2(writeIndex).yx; writeIndex = int2(writeIndex).yx;
} }

2
test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl
View File

@ -4,7 +4,7 @@ using namespace metal;
kernel void f() { kernel void f() {
int const a = 4; int const a = 4;
int3 const b = int3(1, 2, 3); 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; return;
} }

2
test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl
View File

@ -4,7 +4,7 @@ using namespace metal;
kernel void f() { kernel void f() {
int3 const a = int3(1, 2, 3); int3 const a = int3(1, 2, 3);
int const b = 4; 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; return;
} }

2
test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl
View File

@ -4,7 +4,7 @@ using namespace metal;
kernel void f() { kernel void f() {
int3 const a = int3(1, 2, 3); int3 const a = int3(1, 2, 3);
int3 const b = int3(4, 5, 6); 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; return;
} }

2
test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
View File

@ -4,7 +4,7 @@ using namespace metal;
kernel void f() { kernel void f() {
int a = 4; int a = 4;
int3 b = int3(0, 2, 0); 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; return;
} }

2
test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
View File

@ -4,7 +4,7 @@ using namespace metal;
kernel void f() { kernel void f() {
int3 a = int3(1, 2, 3); int3 a = int3(1, 2, 3);
int3 b = int3(0, 5, 0); 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; return;
} }

2
test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
View File

@ -4,7 +4,7 @@ using namespace metal;
kernel void f() { kernel void f() {
int a = 4; int a = 4;
int3 b = int3(0, 2, 0); 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; return;
} }

2
test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
View File

@ -4,7 +4,7 @@ using namespace metal;
kernel void f() { kernel void f() {
int3 a = int3(1, 2, 3); int3 a = int3(1, 2, 3);
int3 b = int3(0, 5, 0); 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; return;
} }

2
test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl
View File

@ -4,7 +4,7 @@ using namespace metal;
kernel void f() { kernel void f() {
int const a = 4; int const a = 4;
int3 const b = int3(1, 2, 3); 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; return;
} }

2
test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl
View File

@ -4,7 +4,7 @@ using namespace metal;
kernel void f() { kernel void f() {
int3 const a = int3(1, 2, 3); int3 const a = int3(1, 2, 3);
int const b = 4; 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; return;
} }

2
test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl
View File

@ -4,7 +4,7 @@ using namespace metal;
kernel void f() { kernel void f() {
int3 const a = int3(1, 2, 3); int3 const a = int3(1, 2, 3);
int3 const b = int3(4, 5, 6); 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; return;
} }

2
test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl
View File

@ -4,7 +4,7 @@ using namespace metal;
kernel void f() { kernel void f() {
int const a = 4; int const a = 4;
int3 const b = int3(1, 2, 3); 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; return;
} }

2
test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl
View File

@ -4,7 +4,7 @@ using namespace metal;
kernel void f() { kernel void f() {
int3 const a = int3(1, 2, 3); int3 const a = int3(1, 2, 3);
int const b = 4; 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; return;
} }

2
test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl
View File

@ -4,7 +4,7 @@ using namespace metal;
kernel void f() { kernel void f() {
int3 const a = int3(1, 2, 3); int3 const a = int3(1, 2, 3);
int3 const b = int3(4, 5, 6); 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; return;
} }

2
test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl
View File

@ -6,6 +6,6 @@ struct S {
}; };
void foo(device S* const tint_symbol) { 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)))));
} }

2
test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl
View File

@ -6,6 +6,6 @@ struct S {
}; };
void foo(device S* const tint_symbol) { 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)))));
} }

2
test/tint/statements/compound_assign/vector/times.wgsl.expected.msl
View File

@ -6,6 +6,6 @@ struct S {
}; };
void foo(device S* const tint_symbol) { 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)))));
} }