From 40141ffddfe6a00fba0cce7e1b5f3759b8d520b8 Mon Sep 17 00:00:00 2001 From: Bill Hollings Date: Wed, 22 Sep 2021 18:58:31 -0400 Subject: MSL: Selectively enable fast-math in MSL code to match Vulkan CTS results. Based on CTS testing, math optimizations between MSL and Vulkan are inconsistent. In some cases, enabling MSL's fast-math compilation option matches Vulkan's math results. In other cases, disabling it does. Broadly enabling or disabling fast-math across all shaders results in some CTS test failures either way. To fix this, selectively enable/disable fast-math optimizations in the MSL code, using metal::fast and metal::precise function namespaces, where supported, and the [[clang::optnone]] function attribute otherwise. Adjust SPIRV-Cross unit test reference shaders to accommodate these changes. --- reference/shaders-msl-no-opt/comp/glsl.std450.comp | 12 ++++++------ reference/shaders-msl-no-opt/frag/fp16.desktop.invalid.frag | 10 +++++----- 2 files changed, 11 insertions(+), 11 deletions(-) (limited to 'reference/shaders-msl-no-opt') diff --git a/reference/shaders-msl-no-opt/comp/glsl.std450.comp b/reference/shaders-msl-no-opt/comp/glsl.std450.comp index 3c505abe..1b3295f4 100644 --- a/reference/shaders-msl-no-opt/comp/glsl.std450.comp +++ b/reference/shaders-msl-no-opt/comp/glsl.std450.comp @@ -170,7 +170,7 @@ float2x2 spvInverse2x2(float2x2 m) } template -inline T spvReflect(T i, T n) +[[clang::optnone]] T spvReflect(T i, T n) { return i - T(2) * i * n * n; } @@ -217,13 +217,13 @@ kernel void main0(device SSBO& _19 [[buffer(0)]]) _19.res = asin(((device float*)&_19.f32)[0u]); _19.res = acos(((device float*)&_19.f32)[0u]); _19.res = atan(((device float*)&_19.f32)[0u]); - _19.res = sinh(((device float*)&_19.f32)[0u]); - _19.res = cosh(((device float*)&_19.f32)[0u]); - _19.res = tanh(((device float*)&_19.f32)[0u]); + _19.res = fast::sinh(((device float*)&_19.f32)[0u]); + _19.res = fast::cosh(((device float*)&_19.f32)[0u]); + _19.res = precise::tanh(((device float*)&_19.f32)[0u]); _19.res = asinh(((device float*)&_19.f32)[0u]); _19.res = acosh(((device float*)&_19.f32)[0u]); _19.res = atanh(((device float*)&_19.f32)[0u]); - _19.res = atan2(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u]); + _19.res = precise::atan2(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u]); _19.res = pow(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u]); _19.res = exp(((device float*)&_19.f32)[0u]); _19.res = log(((device float*)&_19.f32)[0u]); @@ -239,7 +239,7 @@ kernel void main0(device SSBO& _19 [[buffer(0)]]) _19.res = spvRefract(((device float*)&_19.f32)[0u], ((device float*)&_19.f32)[1u], ((device float*)&_19.f32)[2u]); _19.res = length(_19.f32.xy); _19.res = distance(_19.f32.xy, _19.f32.zw); - float2 v2 = normalize(_19.f32.xy); + float2 v2 = fast::normalize(_19.f32.xy); v2 = faceforward(_19.f32.xy, _19.f32.yz, _19.f32.zw); v2 = reflect(_19.f32.xy, _19.f32.zw); v2 = refract(_19.f32.xy, _19.f32.yz, ((device float*)&_19.f32)[3u]); diff --git a/reference/shaders-msl-no-opt/frag/fp16.desktop.invalid.frag b/reference/shaders-msl-no-opt/frag/fp16.desktop.invalid.frag index 3bf42962..de53d681 100644 --- a/reference/shaders-msl-no-opt/frag/fp16.desktop.invalid.frag +++ b/reference/shaders-msl-no-opt/frag/fp16.desktop.invalid.frag @@ -94,11 +94,11 @@ void test_builtins(thread half4& v4, thread half3& v3, thread half& v1) res = cos(v4); res = tan(v4); res = asin(v4); - res = atan2(v4, v3.xyzz); + res = precise::atan2(v4, v3.xyzz); res = atan(v4); - res = sinh(v4); - res = cosh(v4); - res = tanh(v4); + res = fast::sinh(v4); + res = fast::cosh(v4); + res = precise::tanh(v4); res = asinh(v4); res = acosh(v4); res = atanh(v4); @@ -143,7 +143,7 @@ void test_builtins(thread half4& v4, thread half3& v3, thread half& v1) t0 = distance(v4, v4); t0 = dot(v4, v4); half3 res3 = cross(v3, v3); - res = normalize(v4); + res = fast::normalize(v4); res = faceforward(v4, v4, v4); res = reflect(v4, v4); res = refract(v4, v4, v1); -- cgit v1.2.3