Merge pull request #1749 from billhollings/fastmath-quantize

MSL: Honor infinities in OpQuantizeToF16 when compiling using fast-math.

1 files changed, 15 insertions, 4 deletions

diff --git a/reference/opt/shaders-msl/asm/comp/quantize.asm.comp b/reference/opt/shaders-msl/asm/comp/quantize.asm.comp
index 1839ec7a..b7e6f919 100644
--- a/reference/opt/shaders-msl/asm/comp/quantize.asm.comp
+++ b/reference/opt/shaders-msl/asm/comp/quantize.asm.comp
@@ -1,3 +1,5 @@
+#pragma clang diagnostic ignored "-Wmissing-prototypes"
+
 #include <metal_stdlib>
 #include <simd/simd.h>
 
@@ -11,11 +13,20 @@ struct SSBO0
     float4 vec4_val;
 };
 
+template <typename F> struct SpvHalfTypeSelector;
+template <> struct SpvHalfTypeSelector<float> { public: using H = half; };
+template<uint N> struct SpvHalfTypeSelector<vec<float, N>> { using H = vec<half, N>; };
+template<typename F, typename H = typename SpvHalfTypeSelector<F>::H>
+[[clang::optnone]] F spvQuantizeToF16(F val)
+{
+    return F(H(val));
+}
+
 kernel void main0(device SSBO0& _4 [[buffer(0)]])
 {
-    _4.scalar = float(half(_4.scalar));
-    _4.vec2_val = float2(half2(_4.vec2_val));
-    _4.vec3_val = float3(half3(_4.vec3_val));
-    _4.vec4_val = float4(half4(_4.vec4_val));
+    _4.scalar = spvQuantizeToF16(_4.scalar);
+    _4.vec2_val = spvQuantizeToF16(_4.vec2_val);
+    _4.vec3_val = spvQuantizeToF16(_4.vec3_val);
+    _4.vec4_val = spvQuantizeToF16(_4.vec4_val);
 }