GPU: give Metal an IEEE-754 binary64 in software

Metal Shading Language has no double, and the tracking needs one: the track
parametrisation, the propagator and the material budget all do double
arithmetic. The Metal entry point in dev already includes
GPUCommonDoubleBinary64.h; this adds it.

The class implements binary64 over a 64-bit integer. Round to nearest even,
with subnormals, infinities and NaNs, and NaN propagation in the ARM64 order,
so an Apple host is a bit-exact reference down to the payload. Addition,
subtraction, multiplication, division and the conversions to and from float
and the 32-bit integers are exact; sin and cos are fdlibm's and land within
2 ulp of libm. There is no fused multiply-add and no square root. The entry
point aliases the double keyword to the class, so the shared headers go on
saying double, and the header refuses to build where a real double exists.

Abs<double> needs a Metal specialisation of its own, since fabs would
otherwise resolve to metal::fabs(float) through the implicit conversion and
round the mantissa away at 17 call sites, with no error and no warning, among
them the SMatrixGPU pivot selection. A static_assert keeps it that way, since
metal::fabs is not constant-evaluable. SinCosd goes the other way: nothing on
the device path needs it, so math_utils::sincosd is excluded for Metal as it
already is for OpenCL, and the stale GPUCommonDouble.h include goes with it.
The one remaining line is a ternary that needs an explicit value_t to stay
unambiguous once value_t is a class.

It costs of the order of a hundred times plain float on an M-series GPU, which
the tracking can afford because double is a small fraction of its floating
point work.

Author: ktf

Common/MathUtils/include/MathUtils/Utils.h

@@ -111,7 +111,7 @@ GPUdi() void sincos(float ang, float& s, float& c)
 {
   detail::sincos<float>(ang, s, c);
 }
-#ifndef __OPENCL__
+#if !defined(__OPENCL__) && !defined(__METAL__)
 GPUdi() void sincosd(double ang, double& s, double& c)
 {
   detail::sincos<double>(ang, s, c);

DataFormats/Reconstruction/src/TrackParametrizationWithError.cxx

@@ -638,7 +638,7 @@ GPUd() bool TrackParametrizationWithError<value_T>::propagateTo(value_t xk, cons
   }
   double r1pr2Inv = 1. / (r1 + r2), r2inv = 1. / r2, r1inv = 1. / r1;
   double dy2dx = (f1 + f2) * r1pr2Inv, dx2r1pr2 = dx * r1pr2Inv;
-  value_t step = (gpu::CAMath::Abs(x2r) < 0.05f) ? dx * gpu::CAMath::Abs(r2 + f2 * dy2dx)                                                   // chord
+  value_t step = (gpu::CAMath::Abs(x2r) < 0.05f) ? value_t(dx * gpu::CAMath::Abs(r2 + f2 * dy2dx))                                          // chord
                                                  : 2.f * gpu::CAMath::ASin(0.5f * dx * gpu::CAMath::Sqrt(1.f + dy2dx * dy2dx) * crv) / crv; // arc
   step *= gpu::CAMath::Sqrt(1.f + this->getTgl() * this->getTgl());
   //