MSL: Handle implicit integer promotion rules.

MSL inherits the behavior of C where arithmetic on small types are
implicitly converted to int. SPIR-V does not have this behavior, so make
sure that arithmetic results are handled correctly.

2 files changed, 94 insertions, 1 deletions

diff --git a/reference/shaders-msl-no-opt/comp/implicit-integer-promotion.comp b/reference/shaders-msl-no-opt/comp/implicit-integer-promotion.comp
new file mode 100644
index 00000000..5c3ce49e
--- /dev/null
+++ b/reference/shaders-msl-no-opt/comp/implicit-integer-promotion.comp
@@ -0,0 +1,93 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct BUF0
+{
+    half2 f16s;
+    ushort2 u16;
+    short2 i16;
+    ushort4 u16s;
+    short4 i16s;
+    half f16;
+};
+
+static inline __attribute__((always_inline))
+void test_u16(device BUF0& v_24)
+{
+    v_24.f16 += as_type<half>(ushort(((device ushort*)&v_24.u16)[0u] + ((device ushort*)&v_24.u16)[1u]));
+    v_24.f16 += as_type<half>(ushort(((device ushort*)&v_24.u16)[0u] - ((device ushort*)&v_24.u16)[1u]));
+    v_24.f16 += as_type<half>(ushort(((device ushort*)&v_24.u16)[0u] * ((device ushort*)&v_24.u16)[1u]));
+    v_24.f16 += as_type<half>(ushort(((device ushort*)&v_24.u16)[0u] / ((device ushort*)&v_24.u16)[1u]));
+    v_24.f16 += as_type<half>(ushort(((device ushort*)&v_24.u16)[0u] % ((device ushort*)&v_24.u16)[1u]));
+    v_24.f16 += as_type<half>(ushort(((device ushort*)&v_24.u16)[0u] << ((device ushort*)&v_24.u16)[1u]));
+    v_24.f16 += as_type<half>(ushort(((device ushort*)&v_24.u16)[0u] >> ((device ushort*)&v_24.u16)[1u]));
+    v_24.f16 += as_type<half>(ushort(~((device ushort*)&v_24.u16)[0u]));
+    v_24.f16 += as_type<half>(ushort(-((device ushort*)&v_24.u16)[0u]));
+    v_24.f16 += as_type<half>(ushort(((device ushort*)&v_24.u16)[0u] ^ ((device ushort*)&v_24.u16)[1u]));
+    v_24.f16 += as_type<half>(ushort(((device ushort*)&v_24.u16)[0u] & ((device ushort*)&v_24.u16)[1u]));
+    v_24.f16 += as_type<half>(ushort(((device ushort*)&v_24.u16)[0u] | ((device ushort*)&v_24.u16)[1u]));
+}
+
+static inline __attribute__((always_inline))
+void test_i16(device BUF0& v_24)
+{
+    v_24.f16 += as_type<half>(short(((device short*)&v_24.i16)[0u] + ((device short*)&v_24.i16)[1u]));
+    v_24.f16 += as_type<half>(short(((device short*)&v_24.i16)[0u] - ((device short*)&v_24.i16)[1u]));
+    v_24.f16 += as_type<half>(short(((device short*)&v_24.i16)[0u] * ((device short*)&v_24.i16)[1u]));
+    v_24.f16 += as_type<half>(short(((device short*)&v_24.i16)[0u] / ((device short*)&v_24.i16)[1u]));
+    v_24.f16 += as_type<half>(short(((device short*)&v_24.i16)[0u] % ((device short*)&v_24.i16)[1u]));
+    v_24.f16 += as_type<half>(short(((device short*)&v_24.i16)[0u] << ((device short*)&v_24.i16)[1u]));
+    v_24.f16 += as_type<half>(short(((device short*)&v_24.i16)[0u] >> ((device short*)&v_24.i16)[1u]));
+    v_24.f16 += as_type<half>(short(~((device short*)&v_24.i16)[0u]));
+    v_24.f16 += as_type<half>(short(-((device short*)&v_24.i16)[0u]));
+    v_24.f16 += as_type<half>(short(((device short*)&v_24.i16)[0u] ^ ((device short*)&v_24.i16)[1u]));
+    v_24.f16 += as_type<half>(short(((device short*)&v_24.i16)[0u] & ((device short*)&v_24.i16)[1u]));
+    v_24.f16 += as_type<half>(short(((device short*)&v_24.i16)[0u] | ((device short*)&v_24.i16)[1u]));
+}
+
+static inline __attribute__((always_inline))
+void test_u16s(device BUF0& v_24)
+{
+    v_24.f16s += as_type<half2>(v_24.u16s.xy + v_24.u16s.zw);
+    v_24.f16s += as_type<half2>(v_24.u16s.xy - v_24.u16s.zw);
+    v_24.f16s += as_type<half2>(v_24.u16s.xy * v_24.u16s.zw);
+    v_24.f16s += as_type<half2>(v_24.u16s.xy / v_24.u16s.zw);
+    v_24.f16s += as_type<half2>(v_24.u16s.xy % v_24.u16s.zw);
+    v_24.f16s += as_type<half2>(v_24.u16s.xy << v_24.u16s.zw);
+    v_24.f16s += as_type<half2>(v_24.u16s.xy >> v_24.u16s.zw);
+    v_24.f16s += as_type<half2>(~v_24.u16s.xy);
+    v_24.f16s += as_type<half2>(-v_24.u16s.xy);
+    v_24.f16s += as_type<half2>(v_24.u16s.xy ^ v_24.u16s.zw);
+    v_24.f16s += as_type<half2>(v_24.u16s.xy & v_24.u16s.zw);
+    v_24.f16s += as_type<half2>(v_24.u16s.xy | v_24.u16s.zw);
+}
+
+static inline __attribute__((always_inline))
+void test_i16s(device BUF0& v_24)
+{
+    v_24.f16s += as_type<half2>(v_24.i16s.xy + v_24.i16s.zw);
+    v_24.f16s += as_type<half2>(v_24.i16s.xy - v_24.i16s.zw);
+    v_24.f16s += as_type<half2>(v_24.i16s.xy * v_24.i16s.zw);
+    v_24.f16s += as_type<half2>(v_24.i16s.xy / v_24.i16s.zw);
+    v_24.f16s += as_type<half2>(v_24.i16s.xy % v_24.i16s.zw);
+    v_24.f16s += as_type<half2>(v_24.i16s.xy << v_24.i16s.zw);
+    v_24.f16s += as_type<half2>(v_24.i16s.xy >> v_24.i16s.zw);
+    v_24.f16s += as_type<half2>(~v_24.i16s.xy);
+    v_24.f16s += as_type<half2>(-v_24.i16s.xy);
+    v_24.f16s += as_type<half2>(v_24.i16s.xy ^ v_24.i16s.zw);
+    v_24.f16s += as_type<half2>(v_24.i16s.xy & v_24.i16s.zw);
+    v_24.f16s += as_type<half2>(v_24.i16s.xy | v_24.i16s.zw);
+}
+
+kernel void main0(device BUF0& v_24 [[buffer(0)]])
+{
+    test_u16(v_24);
+    test_i16(v_24);
+    test_u16s(v_24);
+    test_i16s(v_24);
+}
+
diff --git a/reference/shaders-msl-no-opt/comp/int16min-literal.comp b/reference/shaders-msl-no-opt/comp/int16min-literal.comp
index a2b36ede..d73768c3 100644
--- a/reference/shaders-msl-no-opt/comp/int16min-literal.comp
+++ b/reference/shaders-msl-no-opt/comp/int16min-literal.comp
@@ -18,7 +18,7 @@ constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
 kernel void main0(constant UBO& _12 [[buffer(0)]], device SSBO& _24 [[buffer(1)]])
 {
     short v = as_type<short>(_12.b);
-    v ^= short(-32768);
+    v = short(v ^ short(-32768));
     _24.a = as_type<half>(v);
 }