Improve handling of INT_MIN/INT64_MIN literals.

We cannot naively convert these to decimal literals. C/C++ (and thus
MSL) has extremely awkward literal promotion rules.

7 files changed, 164 insertions, 0 deletions

diff --git a/reference/shaders-hlsl-no-opt/comp/intmin-literal.comp b/reference/shaders-hlsl-no-opt/comp/intmin-literal.comp
new file mode 100644
index 00000000..9faa7fba
--- /dev/null
+++ b/reference/shaders-hlsl-no-opt/comp/intmin-literal.comp
@@ -0,0 +1,19 @@
+static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u);
+
+RWByteAddressBuffer _9 : register(u1);
+cbuffer UBO : register(b0)
+{
+    float _14_b : packoffset(c0);
+};
+
+
+void comp_main()
+{
+    _9.Store(0, asuint(asfloat(asint(_14_b) ^ int(0x80000000))));
+}
+
+[numthreads(1, 1, 1)]
+void main()
+{
+    comp_main();
+}
diff --git a/reference/shaders-msl-no-opt/comp/int16min-literal.comp b/reference/shaders-msl-no-opt/comp/int16min-literal.comp
new file mode 100644
index 00000000..a2b36ede
--- /dev/null
+++ b/reference/shaders-msl-no-opt/comp/int16min-literal.comp
@@ -0,0 +1,24 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct UBO
+{
+    half b;
+};
+
+struct SSBO
+{
+    half a;
+};
+
+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);
+    _24.a = as_type<half>(v);
+}
+
diff --git a/reference/shaders-msl-no-opt/comp/int64min-literal.msl22.comp b/reference/shaders-msl-no-opt/comp/int64min-literal.msl22.comp
new file mode 100644
index 00000000..a8f2b0e2
--- /dev/null
+++ b/reference/shaders-msl-no-opt/comp/int64min-literal.msl22.comp
@@ -0,0 +1,24 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct UBO
+{
+    float b;
+};
+
+struct SSBO
+{
+    float a;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+kernel void main0(constant UBO& _12 [[buffer(0)]], device SSBO& _25 [[buffer(1)]])
+{
+    long v = long(as_type<int>(_12.b));
+    v ^= long(0x8000000000000000ul);
+    _25.a = as_type<float>(int(v));
+}
+
diff --git a/reference/shaders-msl-no-opt/comp/intmin-literal.comp b/reference/shaders-msl-no-opt/comp/intmin-literal.comp
new file mode 100644
index 00000000..db2294fe
--- /dev/null
+++ b/reference/shaders-msl-no-opt/comp/intmin-literal.comp
@@ -0,0 +1,22 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO
+{
+    float a;
+};
+
+struct UBO
+{
+    float b;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+kernel void main0(device SSBO& _9 [[buffer(0)]], constant UBO& _14 [[buffer(1)]])
+{
+    _9.a = as_type<float>(as_type<int>(_14.b) ^ int(0x80000000));
+}
+
diff --git a/reference/shaders-no-opt/comp/int16min-literal.comp b/reference/shaders-no-opt/comp/int16min-literal.comp
new file mode 100644
index 00000000..d4f3a799
--- /dev/null
+++ b/reference/shaders-no-opt/comp/int16min-literal.comp
@@ -0,0 +1,34 @@
+#version 450
+#if defined(GL_AMD_gpu_shader_int16)
+#extension GL_AMD_gpu_shader_int16 : require
+#elif defined(GL_NV_gpu_shader5)
+#extension GL_NV_gpu_shader5 : require
+#else
+#error No extension available for Int16.
+#endif
+#if defined(GL_AMD_gpu_shader_half_float)
+#extension GL_AMD_gpu_shader_half_float : require
+#elif defined(GL_NV_gpu_shader5)
+#extension GL_NV_gpu_shader5 : require
+#else
+#error No extension available for FP16.
+#endif
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(binding = 0, std140) uniform UBO
+{
+    float16_t b;
+} _12;
+
+layout(binding = 1, std430) buffer SSBO
+{
+    float16_t a;
+} _24;
+
+void main()
+{
+    int16_t v = float16BitsToInt16(_12.b);
+    v ^= (-32768s);
+    _24.a = int16BitsToFloat16(v);
+}
+
diff --git a/reference/shaders-no-opt/comp/int64min-literal.comp b/reference/shaders-no-opt/comp/int64min-literal.comp
new file mode 100644
index 00000000..0eea0d7c
--- /dev/null
+++ b/reference/shaders-no-opt/comp/int64min-literal.comp
@@ -0,0 +1,23 @@
+#version 450
+#extension GL_ARB_gpu_shader_int64 : require
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(binding = 0, std140) uniform UBO
+{
+    float b;
+} _12;
+
+layout(binding = 1, std430) buffer SSBO
+{
+    float a;
+} _32;
+
+void main()
+{
+    double b2 = double(_12.b);
+    int64_t v = doubleBitsToInt64(b2);
+    v ^= int64_t(0x8000000000000000ul);
+    double a2 = int64BitsToDouble(v);
+    _32.a = float(a2);
+}
+
diff --git a/reference/shaders-no-opt/comp/intmin-literal.comp b/reference/shaders-no-opt/comp/intmin-literal.comp
new file mode 100644
index 00000000..5a4896f9
--- /dev/null
+++ b/reference/shaders-no-opt/comp/intmin-literal.comp
@@ -0,0 +1,18 @@
+#version 450
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(binding = 1, std430) buffer SSBO
+{
+    float a;
+} _9;
+
+layout(binding = 0, std140) uniform UBO
+{
+    float b;
+} _14;
+
+void main()
+{
+    _9.a = intBitsToFloat(floatBitsToInt(_14.b) ^ int(0x80000000));
+}
+