Deal with scalar input values for distance/length/normalize.

HLSL and MSL don't support it, so fall back to simpler intrinsics.

reference/opt/shaders-hlsl/comp/scalar-std450-distance-length-normalize.comp

@@ -0,0 +1,14 @@
+RWByteAddressBuffer _9 : register(u0);
+
+void comp_main()
+{
+    _9.Store(8, asuint(distance(asfloat(_9.Load(0)), asfloat(_9.Load(4)))));
+    _9.Store(12, asuint(length(asfloat(_9.Load(0)))));
+    _9.Store(16, asuint(sign(asfloat(_9.Load(0)))));
+}
+
+[numthreads(1, 1, 1)]
+void main()
+{
+    comp_main();
+}

reference/opt/shaders-msl/comp/scalar-std450-distance-length-normalize.comp

@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO
+{
+    float a;
+    float b;
+    float c;
+    float d;
+    float e;
+};
+
+kernel void main0(device SSBO& _9 [[buffer(0)]])
+{
+    _9.c = abs(_9.a - _9.b);
+    _9.d = abs(_9.a);
+    _9.e = sign(_9.a);
+}
+

reference/opt/shaders/comp/scalar-std450-distance-length-normalize.comp

@@ -0,0 +1,19 @@
+#version 450
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(binding = 0, std430) buffer SSBO
+{
+    float a;
+    float b;
+    float c;
+    float d;
+    float e;
+} _9;
+
+void main()
+{
+    _9.c = distance(_9.a, _9.b);
+    _9.d = length(_9.a);
+    _9.e = normalize(_9.a);
+}
+

reference/shaders-hlsl/comp/scalar-std450-distance-length-normalize.comp

@@ -0,0 +1,14 @@
+RWByteAddressBuffer _9 : register(u0);
+
+void comp_main()
+{
+    _9.Store(8, asuint(distance(asfloat(_9.Load(0)), asfloat(_9.Load(4)))));
+    _9.Store(12, asuint(length(asfloat(_9.Load(0)))));
+    _9.Store(16, asuint(sign(asfloat(_9.Load(0)))));
+}
+
+[numthreads(1, 1, 1)]
+void main()
+{
+    comp_main();
+}

reference/shaders-msl/comp/scalar-std450-distance-length-normalize.comp

@@ -0,0 +1,21 @@
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using namespace metal;
+
+struct SSBO
+{
+    float a;
+    float b;
+    float c;
+    float d;
+    float e;
+};
+
+kernel void main0(device SSBO& _9 [[buffer(0)]])
+{
+    _9.c = abs(_9.a - _9.b);
+    _9.d = abs(_9.a);
+    _9.e = sign(_9.a);
+}
+

reference/shaders/comp/scalar-std450-distance-length-normalize.comp

@@ -0,0 +1,19 @@
+#version 450
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+layout(binding = 0, std430) buffer SSBO
+{
+    float a;
+    float b;
+    float c;
+    float d;
+    float e;
+} _9;
+
+void main()
+{
+    _9.c = distance(_9.a, _9.b);
+    _9.d = length(_9.a);
+    _9.e = normalize(_9.a);
+}
+

shaders-hlsl/comp/scalar-std450-distance-length-normalize.comp

@@ -0,0 +1,18 @@
+#version 450
+layout(local_size_x = 1) in;
+
+layout(std430, set = 0, binding = 0) buffer SSBO
+{
+	float a;
+	float b;
+	float c;
+	float d;
+	float e;
+};
+
+void main()
+{
+	c = distance(a, b);
+	d = length(a);
+	e = normalize(a);
+}

shaders-msl/comp/scalar-std450-distance-length-normalize.comp

@@ -0,0 +1,18 @@
+#version 450
+layout(local_size_x = 1) in;
+
+layout(std430, set = 0, binding = 0) buffer SSBO
+{
+	float a;
+	float b;
+	float c;
+	float d;
+	float e;
+};
+
+void main()
+{
+	c = distance(a, b);
+	d = length(a);
+	e = normalize(a);
+}

shaders/comp/scalar-std450-distance-length-normalize.comp

@@ -0,0 +1,18 @@
+#version 450
+layout(local_size_x = 1) in;
+
+layout(std430, set = 0, binding = 0) buffer SSBO
+{
+	float a;
+	float b;
+	float c;
+	float d;
+	float e;
+};
+
+void main()
+{
+	c = distance(a, b);
+	d = length(a);
+	e = normalize(a);
+}

spirv_hlsl.cpp

@@ -3245,6 +3245,17 @@ void CompilerHLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
 		break;
 	}
 
+	case GLSLstd450Normalize:
+		// HLSL does not support scalar versions here.
+		if (expression_type(args[0]).vecsize == 1)
+		{
+			// Returns -1 or 1 for valid input, sign() does the job.
+			emit_unary_func_op(result_type, id, args[0], "sign");
+		}
+		else
+			CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
+		break;
+
 	default:
 		CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
 		break;

spirv_msl.cpp

@@ -4582,6 +4582,43 @@ void CompilerMSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
 		//        GLSLstd450InterpolateAtSample (sample_no_perspective qualifier)
 		//        GLSLstd450InterpolateAtOffset
 
+	case GLSLstd450Distance:
+		// MSL does not support scalar versions here.
+		if (expression_type(args[0]).vecsize == 1)
+		{
+			// Equivalent to length(a - b) -> abs(a - b).
+			emit_op(result_type, id,
+			        join("abs(", to_unpacked_expression(args[0]), " - ", to_unpacked_expression(args[1]), ")"),
+			        should_forward(args[0]) && should_forward(args[1]));
+			inherit_expression_dependencies(id, args[0]);
+			inherit_expression_dependencies(id, args[1]);
+		}
+		else
+			CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
+		break;
+
+	case GLSLstd450Length:
+		// MSL does not support scalar versions here.
+		if (expression_type(args[0]).vecsize == 1)
+		{
+			// Equivalent to abs().
+			emit_unary_func_op(result_type, id, args[0], "abs");
+		}
+		else
+			CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
+		break;
+
+	case GLSLstd450Normalize:
+		// MSL does not support scalar versions here.
+		if (expression_type(args[0]).vecsize == 1)
+		{
+			// Returns -1 or 1 for valid input, sign() does the job.
+			emit_unary_func_op(result_type, id, args[0], "sign");
+		}
+		else
+			CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
+		break;
+
 	default:
 		CompilerGLSL::emit_glsl_op(result_type, id, eop, args, count);
 		break;