Conversation
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.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Your PR AliceO2Group#15866 cannot be merged as is. You should either run
clang-formatyourself and update the pull request, or merge this PR in yours.You can find the AliceO2 coding conventions at https://github.com/AliceO2Group/CodingGuidelines.