[MSL] SPIRV-Cross emits undefined behaviour for OpShiftLeftLogical

[rayanht]

Consider the following shader:

; SPIR-V
; Version: 1.3
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 15
; Schema: 0
               OpCapability Shader
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %1 "main"
               OpExecutionMode %1 LocalSize 1 1 1
               OpDecorate %_struct_2 Block
               OpDecorate %3 DescriptorSet 0
               OpDecorate %3 Binding 0
               OpMemberDecorate %_struct_2 0 Offset 0
       %void = OpTypeVoid
          %5 = OpTypeFunction %void
        %int = OpTypeInt 32 1
  %_struct_2 = OpTypeStruct %int
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
%_ptr_StorageBuffer__struct_2 = OpTypePointer StorageBuffer %_struct_2
      %int_0 = OpConstant %int 0
     %int_n1 = OpConstant %int -1
      %int_2 = OpConstant %int 2
          %3 = OpVariable %_ptr_StorageBuffer__struct_2 StorageBuffer
          %1 = OpFunction %void DontInline %5
         %12 = OpLabel
         %13 = OpAccessChain %_ptr_StorageBuffer_int %3 %int_0
         %14 = OpShiftLeftLogical %int %int_n1 %int_2
               OpStore %13 %14
               OpReturn
               OpFunctionEnd

It is seemingly valid in SPIR-V to apply an OpShiftLeftLogical to a negative base:

spirv-cross emits the following MSL:

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

struct _2
{
    int _m0;
};

kernel void main0(device _2& _3 [[buffer(0)]])
{
    _3._m0 = (-1) << 2;
}

This is an undefined operation as demonstrated by running the Metal validation layers:

interesting_shaders/right_shift_UB_msl.metal:13:19: warning: The result of the left shift is undefined because the left operand is
      negative
    _3._m0 = (-1) << 2;
             ~~~~ ^
1 warning generated.

This is a fuzzer-found bug using SPIRVSmith

[HansKristian-Work]

This seems needlessly nitpicky. Is this actually undefined, or is just pedantic C89 behavior that has been inherited? There are similar cases with right shift arithmetic that works fine on any compiler, but it's still undefined in C (which Metal inherits from).

What would a suggested workaround look like?

[rayanht]

@HansKristian-Work it's actually a bit unclear. I can't find anything about shifting introducing UB in the Metal spec. The Metal validation layers complaining about it might very well be due to behaviour inherited from C.

I don't know what's the course of action here for SPIRV-Cross since the spec and the validation layers seem to diverge. I'd err towards following the spec but I'll also submit a ticket to the Metal bug tracker to clarify this behaviour.