GPU: give Metal an IEEE-754 binary64 in software

Metal Shading Language has no double at all, and the tracking needs one in a
handful of places: the track parametrisation, the propagator and the material
budget all carry double members and do double arithmetic.

GPUCommonDoubleBinary64.h implements binary64 on top of a 64-bit integer. Round
to nearest even only, 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.

GPUCommonMath's own double helpers cannot serve the class: SinCosd's OpenCL arm
passes the cosine by pointer where MSL takes a thread reference, so it does not
compile, and Abs<double> is worse because it does -- fabs resolves to
metal::fabs(float) through the implicit conversion, with no error and no warning,
rounding the mantissa away at 17 call sites including the SMatrixGPU pivot
selection and the Kalman guards against Almost0 and Almost1. Both are defined for
Metal in GPUCommonDouble.h instead, and a static_assert there turns the silent
case into a compile error, since metal::fabs is not constant-evaluable.

The Metal entry point aliases the double keyword to the class, so the shared
headers go on saying double and no call site changes. The header refuses to
build anywhere else, since every other backend has a real double.

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. In exchange a Metal build reproduces the CPU result bit for bit wherever
it uses only the exact operations.

Author: ktf

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());
   //