Merge pull request #865 from KhronosGroup/fix-8632019-02-20

Always value-cast FP16 constants instead of using literals.

15 files changed, 57 insertions, 80 deletions

diff --git a/reference/opt/shaders-msl/comp/bitcast-16bit-1.invalid.comp b/reference/opt/shaders-msl/comp/bitcast-16bit-1.invalid.comp
index f91e5fea..ad9733a8 100644
--- a/reference/opt/shaders-msl/comp/bitcast-16bit-1.invalid.comp
+++ b/reference/opt/shaders-msl/comp/bitcast-16bit-1.invalid.comp
@@ -15,7 +15,7 @@ struct SSBO1
 
 kernel void main0(device SSBO0& _25 [[buffer(0)]], device SSBO1& _39 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
 {
-    _39.outputs[gl_GlobalInvocationID.x].x = int(as_type<uint>(as_type<half2>(_25.inputs[gl_GlobalInvocationID.x].xy) + half2(1.0h)));
+    _39.outputs[gl_GlobalInvocationID.x].x = int(as_type<uint>(as_type<half2>(_25.inputs[gl_GlobalInvocationID.x].xy) + half2(half(1.0))));
     _39.outputs[gl_GlobalInvocationID.x].y = as_type<int>(_25.inputs[gl_GlobalInvocationID.x].zw);
     _39.outputs[gl_GlobalInvocationID.x].z = int(as_type<uint>(ushort2(_25.inputs[gl_GlobalInvocationID.x].xy)));
 }
diff --git a/reference/opt/shaders-msl/frag/16bit-constants.frag b/reference/opt/shaders-msl/frag/16bit-constants.frag
index 2aae1e8a..56c7ea5d 100644
--- a/reference/opt/shaders-msl/frag/16bit-constants.frag
+++ b/reference/opt/shaders-msl/frag/16bit-constants.frag
@@ -13,7 +13,7 @@ struct main0_out
 fragment main0_out main0()
 {
     main0_out out = {};
-    out.foo = 1.0h;
+    out.foo = half(1.0);
     out.bar = 2;
     out.baz = 3u;
     return out;
diff --git a/reference/opt/shaders/comp/bitcast-16bit-1.invalid.comp b/reference/opt/shaders/comp/bitcast-16bit-1.invalid.comp
index a4014e89..501f9795 100644
--- a/reference/opt/shaders/comp/bitcast-16bit-1.invalid.comp
+++ b/reference/opt/shaders/comp/bitcast-16bit-1.invalid.comp
@@ -27,7 +27,7 @@ void main()
 {
     uint ident = gl_GlobalInvocationID.x;
     f16vec2 a = int16BitsToFloat16(_25.inputs[ident].xy);
-    _39.outputs[ident].x = int(packFloat2x16(a + f16vec2(1.0hf)));
+    _39.outputs[ident].x = int(packFloat2x16(a + f16vec2(float16_t(1.0))));
     _39.outputs[ident].y = packInt2x16(_25.inputs[ident].zw);
     _39.outputs[ident].z = int(packUint2x16(u16vec2(_25.inputs[ident].xy)));
 }
diff --git a/reference/opt/shaders/desktop-only/frag/fp16.invalid.desktop.frag b/reference/opt/shaders/desktop-only/frag/fp16.invalid.desktop.frag
index e1fc045c..faf79b2b 100644
--- a/reference/opt/shaders/desktop-only/frag/fp16.invalid.desktop.frag
+++ b/reference/opt/shaders/desktop-only/frag/fp16.invalid.desktop.frag
@@ -30,19 +30,19 @@ f16mat3 test_mat3(f16vec3 a, f16vec3 b, f16vec3 c, f16vec3 d, f16vec3 e, f16vec3
 
 void test_constants()
 {
-    float16_t a = 1.0hf;
-    float16_t b = 1.5hf;
-    float16_t c = -1.5hf;
-    float16_t d = (0.0hf / 0.0hf);
-    float16_t e = (1.0hf / 0.0hf);
-    float16_t f = (-1.0hf / 0.0hf);
-    float16_t g = 1014.0hf;
-    float16_t h = 9.5367431640625e-07hf;
+    float16_t a = float16_t(1.0);
+    float16_t b = float16_t(1.5);
+    float16_t c = float16_t(-1.5);
+    float16_t d = float16_t(0.0 / 0.0);
+    float16_t e = float16_t(1.0 / 0.0);
+    float16_t f = float16_t(-1.0 / 0.0);
+    float16_t g = float16_t(1014.0);
+    float16_t h = float16_t(9.5367431640625e-07);
 }
 
 float16_t test_result()
 {
-    return 1.0hf;
+    return float16_t(1.0);
 }
 
 void test_conversions()
@@ -50,7 +50,7 @@ void test_conversions()
     float16_t one = test_result();
     int a = int(one);
     uint b = uint(one);
-    bool c = one != 0.0hf;
+    bool c = one != float16_t(0.0);
     float d = float(one);
     double e = double(one);
     float16_t a2 = float16_t(a);
diff --git a/reference/opt/shaders/frag/16bit-constants.frag b/reference/opt/shaders/frag/16bit-constants.frag
index f9ce356f..57d82561 100644
--- a/reference/opt/shaders/frag/16bit-constants.frag
+++ b/reference/opt/shaders/frag/16bit-constants.frag
@@ -18,7 +18,7 @@ layout(location = 2) out uint16_t baz;
 
 void main()
 {
-    foo = 1.0hf;
+    foo = float16_t(1.0);
     bar = 2s;
     baz = 3us;
 }
diff --git a/reference/shaders-msl/comp/bitcast-16bit-1.invalid.comp b/reference/shaders-msl/comp/bitcast-16bit-1.invalid.comp
index 8315f792..170e4920 100644
--- a/reference/shaders-msl/comp/bitcast-16bit-1.invalid.comp
+++ b/reference/shaders-msl/comp/bitcast-16bit-1.invalid.comp
@@ -17,7 +17,7 @@ kernel void main0(device SSBO0& _25 [[buffer(0)]], device SSBO1& _39 [[buffer(1)
 {
     uint ident = gl_GlobalInvocationID.x;
     half2 a = as_type<half2>(_25.inputs[ident].xy);
-    _39.outputs[ident].x = int(as_type<uint>(a + half2(1.0h)));
+    _39.outputs[ident].x = int(as_type<uint>(a + half2(half(1.0))));
     _39.outputs[ident].y = as_type<int>(_25.inputs[ident].zw);
     _39.outputs[ident].z = int(as_type<uint>(ushort2(_25.inputs[ident].xy)));
 }
diff --git a/reference/shaders-msl/frag/16bit-constants.frag b/reference/shaders-msl/frag/16bit-constants.frag
index 2aae1e8a..56c7ea5d 100644
--- a/reference/shaders-msl/frag/16bit-constants.frag
+++ b/reference/shaders-msl/frag/16bit-constants.frag
@@ -13,7 +13,7 @@ struct main0_out
 fragment main0_out main0()
 {
     main0_out out = {};
-    out.foo = 1.0h;
+    out.foo = half(1.0);
     out.bar = 2;
     out.baz = 3u;
     return out;
diff --git a/reference/shaders-msl/frag/fp16.desktop.invalid.frag b/reference/shaders-msl/frag/fp16.desktop.invalid.frag
index b914e4b8..001944fc 100644
--- a/reference/shaders-msl/frag/fp16.desktop.invalid.frag
+++ b/reference/shaders-msl/frag/fp16.desktop.invalid.frag
@@ -52,19 +52,19 @@ half3x3 test_mat3(thread const half3& a, thread const half3& b, thread const hal
 
 void test_constants()
 {
-    half a = 1.0h;
-    half b = 1.5h;
-    half c = -1.5h;
-    half d = (0.0h / 0.0h);
-    half e = (1.0h / 0.0h);
-    half f = (-1.0h / 0.0h);
-    half g = 1014.0h;
-    half h = 9.5367431640625e-07h;
+    half a = half(1.0);
+    half b = half(1.5);
+    half c = half(-1.5);
+    half d = half(0.0 / 0.0);
+    half e = half(1.0 / 0.0);
+    half f = half(-1.0 / 0.0);
+    half g = half(1014.0);
+    half h = half(9.5367431640625e-07);
 }
 
 half test_result()
 {
-    return 1.0h;
+    return half(1.0);
 }
 
 void test_conversions()
@@ -72,7 +72,7 @@ void test_conversions()
     half one = test_result();
     int a = int(one);
     uint b = uint(one);
-    bool c = one != 0.0h;
+    bool c = one != half(0.0);
     float d = float(one);
     half a2 = half(a);
     half b2 = half(b);
diff --git a/reference/shaders/comp/bitcast-16bit-1.invalid.comp b/reference/shaders/comp/bitcast-16bit-1.invalid.comp
index a4014e89..501f9795 100644
--- a/reference/shaders/comp/bitcast-16bit-1.invalid.comp
+++ b/reference/shaders/comp/bitcast-16bit-1.invalid.comp
@@ -27,7 +27,7 @@ void main()
 {
     uint ident = gl_GlobalInvocationID.x;
     f16vec2 a = int16BitsToFloat16(_25.inputs[ident].xy);
-    _39.outputs[ident].x = int(packFloat2x16(a + f16vec2(1.0hf)));
+    _39.outputs[ident].x = int(packFloat2x16(a + f16vec2(float16_t(1.0))));
     _39.outputs[ident].y = packInt2x16(_25.inputs[ident].zw);
     _39.outputs[ident].z = int(packUint2x16(u16vec2(_25.inputs[ident].xy)));
 }
diff --git a/reference/shaders/desktop-only/frag/fp16.invalid.desktop.frag b/reference/shaders/desktop-only/frag/fp16.invalid.desktop.frag
index e1fc045c..faf79b2b 100644
--- a/reference/shaders/desktop-only/frag/fp16.invalid.desktop.frag
+++ b/reference/shaders/desktop-only/frag/fp16.invalid.desktop.frag
@@ -30,19 +30,19 @@ f16mat3 test_mat3(f16vec3 a, f16vec3 b, f16vec3 c, f16vec3 d, f16vec3 e, f16vec3
 
 void test_constants()
 {
-    float16_t a = 1.0hf;
-    float16_t b = 1.5hf;
-    float16_t c = -1.5hf;
-    float16_t d = (0.0hf / 0.0hf);
-    float16_t e = (1.0hf / 0.0hf);
-    float16_t f = (-1.0hf / 0.0hf);
-    float16_t g = 1014.0hf;
-    float16_t h = 9.5367431640625e-07hf;
+    float16_t a = float16_t(1.0);
+    float16_t b = float16_t(1.5);
+    float16_t c = float16_t(-1.5);
+    float16_t d = float16_t(0.0 / 0.0);
+    float16_t e = float16_t(1.0 / 0.0);
+    float16_t f = float16_t(-1.0 / 0.0);
+    float16_t g = float16_t(1014.0);
+    float16_t h = float16_t(9.5367431640625e-07);
 }
 
 float16_t test_result()
 {
-    return 1.0hf;
+    return float16_t(1.0);
 }
 
 void test_conversions()
@@ -50,7 +50,7 @@ void test_conversions()
     float16_t one = test_result();
     int a = int(one);
     uint b = uint(one);
-    bool c = one != 0.0hf;
+    bool c = one != float16_t(0.0);
     float d = float(one);
     double e = double(one);
     float16_t a2 = float16_t(a);
diff --git a/reference/shaders/frag/16bit-constants.frag b/reference/shaders/frag/16bit-constants.frag
index f9ce356f..57d82561 100644
--- a/reference/shaders/frag/16bit-constants.frag
+++ b/reference/shaders/frag/16bit-constants.frag
@@ -18,7 +18,7 @@ layout(location = 2) out uint16_t baz;
 
 void main()
 {
-    foo = 1.0hf;
+    foo = float16_t(1.0);
     bar = 2s;
     baz = 3us;
 }
diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp
index 74724e79..affaf697 100644
--- a/spirv_glsl.cpp
+++ b/spirv_glsl.cpp
@@ -2961,51 +2961,31 @@ string CompilerGLSL::convert_half_to_string(const SPIRConstant &c, uint32_t col,
 	string res;
 	float float_value = c.scalar_f16(col, row);
 
+	// There is no literal "hf" in GL_NV_gpu_shader5, so to avoid lots
+	// of complicated workarounds, just value-cast to the half type always.
 	if (std::isnan(float_value) || std::isinf(float_value))
 	{
-		if (backend.half_literal_suffix)
-		{
-			// There is no uintBitsToFloat for 16-bit, so have to rely on legacy fallback here.
-			if (float_value == numeric_limits<float>::infinity())
-				res = join("(1.0", backend.half_literal_suffix, " / 0.0", backend.half_literal_suffix, ")");
-			else if (float_value == -numeric_limits<float>::infinity())
-				res = join("(-1.0", backend.half_literal_suffix, " / 0.0", backend.half_literal_suffix, ")");
-			else if (std::isnan(float_value))
-				res = join("(0.0", backend.half_literal_suffix, " / 0.0", backend.half_literal_suffix, ")");
-			else
-				SPIRV_CROSS_THROW("Cannot represent non-finite floating point constant.");
-		}
+		SPIRType type;
+		type.basetype = SPIRType::Half;
+		type.vecsize = 1;
+		type.columns = 1;
+
+		if (float_value == numeric_limits<float>::infinity())
+			res = join(type_to_glsl(type), "(1.0 / 0.0)");
+		else if (float_value == -numeric_limits<float>::infinity())
+			res = join(type_to_glsl(type), "(-1.0 / 0.0)");
+		else if (std::isnan(float_value))
+			res = join(type_to_glsl(type), "(0.0 / 0.0)");
 		else
-		{
-			SPIRType type;
-			type.basetype = SPIRType::Half;
-			type.vecsize = 1;
-			type.columns = 1;
-
-			if (float_value == numeric_limits<float>::infinity())
-				res = join(type_to_glsl(type), "(1.0 / 0.0)");
-			else if (float_value == -numeric_limits<float>::infinity())
-				res = join(type_to_glsl(type), "(-1.0 / 0.0)");
-			else if (std::isnan(float_value))
-				res = join(type_to_glsl(type), "(0.0 / 0.0)");
-			else
-				SPIRV_CROSS_THROW("Cannot represent non-finite floating point constant.");
-		}
+			SPIRV_CROSS_THROW("Cannot represent non-finite floating point constant.");
 	}
 	else
 	{
-		if (backend.half_literal_suffix)
-			res = convert_to_string(float_value) + backend.half_literal_suffix;
-		else
-		{
-			// In HLSL (FXC), it's important to cast the literals to half precision right away.
-			// There is no literal for it.
-			SPIRType type;
-			type.basetype = SPIRType::Half;
-			type.vecsize = 1;
-			type.columns = 1;
-			res = join(type_to_glsl(type), "(", convert_to_string(float_value), ")");
-		}
+		SPIRType type;
+		type.basetype = SPIRType::Half;
+		type.vecsize = 1;
+		type.columns = 1;
+		res = join(type_to_glsl(type), "(", convert_to_string(float_value), ")");
 	}
 
 	return res;
diff --git a/spirv_glsl.hpp b/spirv_glsl.hpp
index f94043b9..7477cf3b 100644
--- a/spirv_glsl.hpp
+++ b/spirv_glsl.hpp
@@ -380,7 +380,6 @@ protected:
 		const char *basic_uint8_type = "uint8_t";
 		const char *basic_int16_type = "int16_t";
 		const char *basic_uint16_type = "uint16_t";
-		const char *half_literal_suffix = "hf";
 		const char *int16_t_literal_suffix = "s";
 		const char *uint16_t_literal_suffix = "us";
 		bool swizzle_is_function = false;
diff --git a/spirv_hlsl.cpp b/spirv_hlsl.cpp
index c3191577..1b9a8270 100644
--- a/spirv_hlsl.cpp
+++ b/spirv_hlsl.cpp
@@ -4618,7 +4618,6 @@ string CompilerHLSL::compile()
 	options.vulkan_semantics = true;
 	backend.float_literal_suffix = true;
 	backend.double_literal_suffix = false;
-	backend.half_literal_suffix = nullptr;
 	backend.long_long_literal_suffix = true;
 	backend.uint32_t_literal_suffix = true;
 	backend.int16_t_literal_suffix = nullptr;
diff --git a/spirv_msl.cpp b/spirv_msl.cpp
index 1d369469..2682ae0a 100644
--- a/spirv_msl.cpp
+++ b/spirv_msl.cpp
@@ -564,7 +564,6 @@ string CompilerMSL::compile()
 	options.version = 450;
 	backend.null_pointer_literal = "nullptr";
 	backend.float_literal_suffix = false;
-	backend.half_literal_suffix = "h";
 	backend.uint32_t_literal_suffix = true;
 	backend.int16_t_literal_suffix = nullptr;
 	backend.uint16_t_literal_suffix = "u";