Merge pull request #1757 from KhronosGroup/fix-1754

Improve handling of INT_MIN/INT64_MIN literals.

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();
+}

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);
+}
+

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));
+}
+

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));
+}
+

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);
+}
+

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);
+}
+

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));
+}
+

shaders-hlsl-no-opt/comp/intmin-literal.comp

@@ -0,0 +1,18 @@
+#version 450
+
+layout(local_size_x = 1) in;
+
+layout(set = 0, binding = 1) buffer SSBO
+{
+	float a;
+};
+
+layout(set = 0, binding = 0) uniform UBO
+{
+	float b;
+};
+
+void main()
+{
+	a = intBitsToFloat(floatBitsToInt(b) ^ 0x80000000);
+}

shaders-msl-no-opt/comp/int16min-literal.comp

@@ -0,0 +1,22 @@
+#version 450
+#extension GL_EXT_shader_explicit_arithmetic_types_int16 : require
+#extension GL_EXT_shader_explicit_arithmetic_types_float16 : require
+
+layout(local_size_x = 1) in;
+
+layout(set = 0, binding = 1) buffer SSBO
+{
+	float16_t a;
+};
+
+layout(set = 0, binding = 0) uniform UBO
+{
+	float16_t b;
+};
+
+void main()
+{
+	int16_t v = float16BitsToInt16(b);
+	v ^= 0x8000s;
+	a = int16BitsToFloat16(v);
+}

shaders-msl-no-opt/comp/int64min-literal.msl22.comp

@@ -0,0 +1,21 @@
+#version 450
+#extension GL_ARB_gpu_shader_int64 : require
+
+layout(local_size_x = 1) in;
+
+layout(set = 0, binding = 1) buffer SSBO
+{
+	float a;
+};
+
+layout(set = 0, binding = 0) uniform UBO
+{
+	float b;
+};
+
+void main()
+{
+	int64_t v = int64_t(floatBitsToInt(b));
+	v ^= 0x8000000000000000L;
+	a = intBitsToFloat(int(v));
+}

shaders-msl-no-opt/comp/intmin-literal.comp

@@ -0,0 +1,18 @@
+#version 450
+
+layout(local_size_x = 1) in;
+
+layout(set = 0, binding = 1) buffer SSBO
+{
+	float a;
+};
+
+layout(set = 0, binding = 0) uniform UBO
+{
+	float b;
+};
+
+void main()
+{
+	a = intBitsToFloat(floatBitsToInt(b) ^ 0x80000000);
+}

shaders-no-opt/comp/int16min-literal.comp

@@ -0,0 +1,22 @@
+#version 450
+#extension GL_EXT_shader_explicit_arithmetic_types_int16 : require
+#extension GL_EXT_shader_explicit_arithmetic_types_float16 : require
+
+layout(local_size_x = 1) in;
+
+layout(set = 0, binding = 1) buffer SSBO
+{
+	float16_t a;
+};
+
+layout(set = 0, binding = 0) uniform UBO
+{
+	float16_t b;
+};
+
+void main()
+{
+	int16_t v = float16BitsToInt16(b);
+	v ^= 0x8000s;
+	a = int16BitsToFloat16(v);
+}

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) in;
+
+layout(set = 0, binding = 1) buffer SSBO
+{
+	float a;
+};
+
+layout(set = 0, binding = 0) uniform UBO
+{
+	float b;
+};
+
+void main()
+{
+	double b2 = b;
+	int64_t v = doubleBitsToInt64(b2);
+	v ^= 0x8000000000000000L;
+	double a2 = int64BitsToDouble(v);
+	a = float(a2);
+}

shaders-no-opt/comp/intmin-literal.comp

@@ -0,0 +1,18 @@
+#version 450
+
+layout(local_size_x = 1) in;
+
+layout(set = 0, binding = 1) buffer SSBO
+{
+	float a;
+};
+
+layout(set = 0, binding = 0) uniform UBO
+{
+	float b;
+};
+
+void main()
+{
+	a = intBitsToFloat(floatBitsToInt(b) ^ 0x80000000);
+}

spirv_common.hpp

@@ -211,6 +211,28 @@ inline std::string convert_to_string(const T &t)
 	return std::to_string(t);
 }
 
+static inline std::string convert_to_string(int32_t value)
+{
+	// INT_MIN is ... special on some backends. If we use a decimal literal, and negate it, we
+	// could accidentally promote the literal to long first, then negate.
+	// To workaround it, emit int(0x80000000) instead.
+	if (value == std::numeric_limits<int32_t>::min())
+		return "int(0x80000000)";
+	else
+		return std::to_string(value);
+}
+
+static inline std::string convert_to_string(int64_t value, const std::string &int64_type, bool long_long_literal_suffix)
+{
+	// INT64_MIN is ... special on some backends.
+	// If we use a decimal literal, and negate it, we might overflow the representable numbers.
+	// To workaround it, emit int(0x80000000) instead.
+	if (value == std::numeric_limits<int64_t>::min())
+		return join(int64_type, "(0x8000000000000000u", (long_long_literal_suffix ? "ll" : "l"), ")");
+	else
+		return std::to_string(value) + (long_long_literal_suffix ? "ll" : "l");
+}
+
 // Allow implementations to set a convenient standard precision
 #ifndef SPIRV_CROSS_FLT_FMT
 #define SPIRV_CROSS_FLT_FMT "%.32g"

spirv_glsl.cpp

@@ -5290,13 +5290,15 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
 		break;
 
 	case SPIRType::Int64:
+	{
+		auto tmp = type;
+		tmp.vecsize = 1;
+		tmp.columns = 1;
+		auto int64_type = type_to_glsl(tmp);
+
 		if (splat)
 		{
-			res += convert_to_string(c.scalar_i64(vector, 0));
+			res += convert_to_string(c.scalar_i64(vector, 0), int64_type, backend.long_long_literal_suffix);
-			if (backend.long_long_literal_suffix)
-				res += "ll";
-			else
-				res += "l";
 		}
 		else
 		{
@@ -5305,19 +5307,14 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t
 				if (c.vector_size() > 1 && c.specialization_constant_id(vector, i) != 0)
 					res += to_name(c.specialization_constant_id(vector, i));
 				else
-				{
+					res += convert_to_string(c.scalar_i64(vector, i), int64_type, backend.long_long_literal_suffix);
-					res += convert_to_string(c.scalar_i64(vector, i));
-					if (backend.long_long_literal_suffix)
-						res += "ll";
-					else
-						res += "l";
-				}
 
 				if (i + 1 < c.vector_size())
 					res += ", ";
 			}
 		}
 		break;
+	}
 
 	case SPIRType::UInt64:
 		if (splat)