diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc index 792b5b5a35..27a51919d2 100644 --- a/src/tint/writer/msl/generator_impl.cc +++ b/src/tint/writer/msl/generator_impl.cc @@ -590,8 +590,18 @@ bool GeneratorImpl::EmitBinary(std::ostream& out, const ast::BinaryExpression* e ScopedParen sp(out); { ScopedBitCast lhs_uint_cast(this, out, lhs_type, unsigned_type_of(lhs_type)); - if (!EmitExpression(out, expr->lhs)) { - return false; + + // 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()) { + ScopedCast lhs_self_cast(this, out, lhs_type, lhs_type); + if (!EmitExpression(out, expr->lhs)) { + return false; + } + } else { + if (!EmitExpression(out, expr->lhs)) { + return false; + } } } if (!emit_op()) { diff --git a/test/tint/bug/tint/1542.wgsl b/test/tint/bug/tint/1542.wgsl new file mode 100644 index 0000000000..32aba351ae --- /dev/null +++ b/test/tint/bug/tint/1542.wgsl @@ -0,0 +1,11 @@ +struct UniformBuffer { + d: vec3, +} + +@group(0) @binding(0) +var u_input: UniformBuffer; + +@compute @workgroup_size(1) +fn main() { + let temp: vec3 = (u_input.d << vec3()); +} diff --git a/test/tint/bug/tint/1542.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1542.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..32c3b0981d --- /dev/null +++ b/test/tint/bug/tint/1542.wgsl.expected.dxc.hlsl @@ -0,0 +1,9 @@ +cbuffer cbuffer_u_input : register(b0, space0) { + uint4 u_input[1]; +}; + +[numthreads(1, 1, 1)] +void main() { + const int3 temp = (asint(u_input[0].xyz) << (0u).xxx); + return; +} diff --git a/test/tint/bug/tint/1542.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1542.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..32c3b0981d --- /dev/null +++ b/test/tint/bug/tint/1542.wgsl.expected.fxc.hlsl @@ -0,0 +1,9 @@ +cbuffer cbuffer_u_input : register(b0, space0) { + uint4 u_input[1]; +}; + +[numthreads(1, 1, 1)] +void main() { + const int3 temp = (asint(u_input[0].xyz) << (0u).xxx); + return; +} diff --git a/test/tint/bug/tint/1542.wgsl.expected.glsl b/test/tint/bug/tint/1542.wgsl.expected.glsl new file mode 100644 index 0000000000..f49dc3bb48 --- /dev/null +++ b/test/tint/bug/tint/1542.wgsl.expected.glsl @@ -0,0 +1,16 @@ +#version 310 es + +layout(binding = 0, std140) uniform UniformBuffer_ubo { + ivec3 d; + uint pad; +} u_input; + +void tint_symbol() { + ivec3 temp = (u_input.d << uvec3(0u)); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + tint_symbol(); + return; +} diff --git a/test/tint/bug/tint/1542.wgsl.expected.msl b/test/tint/bug/tint/1542.wgsl.expected.msl new file mode 100644 index 0000000000..7d8ffdd8be --- /dev/null +++ b/test/tint/bug/tint/1542.wgsl.expected.msl @@ -0,0 +1,26 @@ +#include + +using namespace metal; + +template +struct tint_array { + const constant T& operator[](size_t i) const constant { return elements[i]; } + device T& operator[](size_t i) device { return elements[i]; } + const device T& operator[](size_t i) const device { return elements[i]; } + thread T& operator[](size_t i) thread { return elements[i]; } + const thread T& operator[](size_t i) const thread { return elements[i]; } + threadgroup T& operator[](size_t i) threadgroup { return elements[i]; } + const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; } + T elements[N]; +}; + +struct UniformBuffer { + /* 0x0000 */ packed_int3 d; + /* 0x000c */ tint_array tint_pad; +}; + +kernel void tint_symbol(const constant UniformBuffer* tint_symbol_1 [[buffer(0)]]) { + int3 const temp = as_type((as_type(int3((*(tint_symbol_1)).d)) << uint3(0u))); + return; +} + diff --git a/test/tint/bug/tint/1542.wgsl.expected.spvasm b/test/tint/bug/tint/1542.wgsl.expected.spvasm new file mode 100644 index 0000000000..a66b08a55b --- /dev/null +++ b/test/tint/bug/tint/1542.wgsl.expected.spvasm @@ -0,0 +1,37 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 18 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %UniformBuffer "UniformBuffer" + OpMemberName %UniformBuffer 0 "d" + OpName %u_input "u_input" + OpName %main "main" + OpDecorate %UniformBuffer Block + OpMemberDecorate %UniformBuffer 0 Offset 0 + OpDecorate %u_input NonWritable + OpDecorate %u_input DescriptorSet 0 + OpDecorate %u_input Binding 0 + %int = OpTypeInt 32 1 + %v3int = OpTypeVector %int 3 +%UniformBuffer = OpTypeStruct %v3int +%_ptr_Uniform_UniformBuffer = OpTypePointer Uniform %UniformBuffer + %u_input = OpVariable %_ptr_Uniform_UniformBuffer Uniform + %void = OpTypeVoid + %6 = OpTypeFunction %void + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_v3int = OpTypePointer Uniform %v3int + %v3uint = OpTypeVector %uint 3 + %16 = OpConstantNull %v3uint + %main = OpFunction %void None %6 + %9 = OpLabel + %13 = OpAccessChain %_ptr_Uniform_v3int %u_input %uint_0 + %14 = OpLoad %v3int %13 + %17 = OpShiftLeftLogical %v3int %14 %16 + OpReturn + OpFunctionEnd diff --git a/test/tint/bug/tint/1542.wgsl.expected.wgsl b/test/tint/bug/tint/1542.wgsl.expected.wgsl new file mode 100644 index 0000000000..ca7737e340 --- /dev/null +++ b/test/tint/bug/tint/1542.wgsl.expected.wgsl @@ -0,0 +1,10 @@ +struct UniformBuffer { + d : vec3, +} + +@group(0) @binding(0) var u_input : UniformBuffer; + +@compute @workgroup_size(1) +fn main() { + let temp : vec3 = (u_input.d << vec3()); +} diff --git a/test/tint/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl b/test/tint/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl index f97db25d1b..1e4e146d2f 100644 --- a/test/tint/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl +++ b/test/tint/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl @@ -4,7 +4,7 @@ using namespace metal; kernel void f() { int3 const a = int3(1, 2, 3); uint3 const b = uint3(4u, 5u, 6u); - int3 const r = as_type((as_type(a) << b)); + int3 const r = as_type((as_type(int3(a)) << b)); return; } diff --git a/test/tint/statements/compound_assign/vector/shift_left.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/shift_left.wgsl.expected.msl index 2dd819cebe..84a60be660 100644 --- a/test/tint/statements/compound_assign/vector/shift_left.wgsl.expected.msl +++ b/test/tint/statements/compound_assign/vector/shift_left.wgsl.expected.msl @@ -6,6 +6,6 @@ struct S { }; void foo(device S* const tint_symbol) { - (*(tint_symbol)).a = as_type((as_type((*(tint_symbol)).a) << uint4(2u))); + (*(tint_symbol)).a = as_type((as_type(int4((*(tint_symbol)).a)) << uint4(2u))); }