Conversation
Member
Author
|
@davidrohr this actually seems to compile fine. I have yet to review all the commits. There is a bunch of them which are to fix all the places which need As discussed privately, one nice side effect of this development is that due to the fact metal does not support double, there is now a path which uses two floats per double to carry around better precision than merely doing all the calculations as float. If confirmed, the synthetic propagation benchmarks using it are actually quite nice, at least on my M1 Max. I will bug you next to see how to setup a proper benchmark, if you are around. |
Collaborator
Collaborator
|
Disentagling the GPUglobalconstexpr makes sense in my opinion. Although, I would file a bug report to metal asking why they do not support simply constexpr. I do not see any reason not to, and apparently all other GPU APIs support it, but then Apple is Apple... |
macOS ships OpenCL 1.2, below the 2.x the OpenCL backend requires, so find_package(OpenCL) there could never produce a usable backend: the version check dropped it again a few lines later. Skip the lookup on Apple instead. With that, CUDA_ENABLED, OPENCL_ENABLED and HIP_ENABLED are all necessarily off on macOS, which makes the Darwin arm of the backend dispatch dead code. Drop it along with its warning and unindent the rest.
The MSL generic address space covers device, threadgroup and thread but not constant, so a generic member function cannot be called on an object that lives in constant memory: 'cannot initialize object parameter of type X with an expression of type constant X'. The shared code is generic throughout, so constant memory is simply not usable on this backend. Metal therefore implies GPUCA_NO_CONSTANT_MEMORY, which already exists for the other backends and redirects GPUconstant() to GPUglobal(). GPUconstantref() has to follow it, exactly as the OpenCL block already arranges; the Metal block hardcoded 'constant' and so kept handing out constant references whatever the setting. It now falls through to the unannotated, and therefore generic, fallback. Macro expansions are unchanged for host, CUDA, HIP, OpenCL and cling. Takes the Metal translation unit from 333 errors to 235, of which the constant-versus-generic diagnostics drop from 91 to 14.
The backend itself: the Objective-C++ host side, the .metal kernel source and its build rules, plus the CMake to enable them. Off unless asked for. FindO2GPU.cmake leaves ENABLE_METAL=OFF on macOS and the subdirectory is gated on METAL_ENABLED, so macOS keeps running on the CPU until the whole chain is validated. Apple toolchain only: the source goes .metal -> AIR through xcrun metal and nothing else, with no SPIR-V translation step in between. Requires -std=metal4.1, the first MSL version with a generic address space. Earlier versions reject an unannotated pointer with 'pointer type must have explicit address space qualifier' and an unannotated 'this' with 'cannot initialize object parameter', both of which GPUCommonDefAPI.h relies on for GPUgeneric() and GPUdDefault(). Verified against Xcode 27, which ships metal4.1; Xcode 26 and earlier stop at metal4.0. Two things that look structural are not. MSL rejects derived classes, but '#pragma METAL internals : enable' -- the switch metal_stdlib itself uses in 46 paired places -- lifts that, and the kernels are class-based throughout. A derived type still cannot be a kernel argument, so the constant memory arrives as an untyped buffer and is cast inside, mirroring what the OpenCL TU does with __cl_clang_non_portable_kernel_param_types and what gpu_mem already did here. That also settles the constant address space, which generic does not span. With both in place the kernel list expands to all 104 entry points and no derived-class or kernel-argument-type errors remain. The bodies still do not compile: 1038 errors, of which 360 are MSL having no double (largely host-only headers such as PhysicsConstants.h and MathUtils/Utils.h reaching device code) and 352 are namespace-scope constexpr needing GPUglobalconstexpr(). That is bulk work rather than a missing language feature, so the .metal file still stops after the common headers until it is done.
MSL has no double, and rejects it at the declaration, so PadPlane could not be declared at all in a Metal translation unit -- never mind read. It is embedded by value in GeometryBase, which is embedded in GeometryFlat, which GPUTRDGeometry inherits, so this blocked the whole TRD geometry on the device. Storage is deliberately left alone. GPUdoubleStore occupies the same eight bytes with the same alignment as a double, so the object layout is identical on both sides and the host still writes and uploads exactly the bytes it did before; only the device declaration differs, and the stored bits are decoded to float on read. On every other backend GPUdoubleStore is a plain double and GPUdoubleGet() is the identity, so CUDA, HIP and the host are untouched -- including ROOT I/O, which only ever sees the host type. Narrowing to float on the device costs nothing that was not already lost: GPUTRDGeometry truncates every one of these accessors to float anyway. The decoder is exact, not approximate: round to nearest even, with the subnormal, infinity and NaN cases handled. It agrees bit-for-bit with a real double to float conversion over four million values, two million uniform across the geometry range and two million random bit patterns. A faster branchless variant is possible and is left for later if it ever shows up in a profile. Takes PadPlane.h from 73 errors to 0 in a Metal translation unit.
…ut double Two headers reach GPUConstantMem.h through PID.h and Propagator.h and account for most of the remaining double in device code. Neither needs the PadPlane treatment, because neither holds transferred state: these are compile-time constants and free functions, so there is no layout to preserve. PhysicsConstants.h: the particle masses only ever serve as compile-time initialisers, and the table built from them, PID::sMasses, is float already. They now use a MassType alias that is float on Metal and double everywhere else, so nothing is lost that was not already being narrowed. The declarations sit in a generated block, so make_pdg_header.py emits the same thing and a regeneration will not undo this. Checked that the double to float narrowing is exact for all 123 literal masses: none of them hits a double rounding case, so sMasses comes out bit identical either way. MathUtils/Utils.h: every helper here has a float version and a d suffixed double twin, and nothing under GPU/ calls a d variant. The double twins are now compiled out on Metal only. Guarded on __METAL__ rather than GPUCA_GPUCODE_DEVICE on purpose. The latter is also set for the CUDA, HIP and OpenCL device passes, which do have double, and using it would have silently dropped those helpers from CUDA and turned its masses into floats. Preprocessed declarations are unchanged for host, CUDA, HIP, OpenCL and cling. Together this takes both headers from 226 errors to 0 in a Metal translation unit, and the whole kernel translation unit from 820 to 590.
TrackParametrizationWithError declares MatrixDSym5 and MatrixD5 as SMatrix over double. MSL has no double, so those two alias declarations fail, and because the failure is mid-declaration the parser then loses the rest of the alias and every later mention of the name. The class comes out malformed, which is why TrackParCov stopped being recognised as a base of TrackTPCITS and why every subsequent call to a TrackParametrization method through it reported 'cannot initialize object parameter ... with an expression of type const TrackParCov'. Those diagnostics read like an address-space problem and are not one: they are all fallout from these two lines. Guarding the aliases, and the three declarations that mention them, removes all of them. The three methods are declared here but defined out of line, so device code could not call them on Metal in any case. Preprocessed output is unchanged for host, CUDA, HIP and cling; the diff adds only #ifndef __METAL__ guards and removes nothing. Takes the Metal translation unit from 235 errors to 199, with the constant-versus-generic diagnostics down from 14 to 0.
Two things stop the GPU SMatrix port from being usable on Metal, neither of them to do with the matrix maths. MSL rejects variables declared static at function scope, which is how MatRepSymGPU::off() caches its offset table. Function-scope constexpr without static is accepted, so Metal uses that; every other backend keeps the static. The streaming operator is constrained with a C++20 requires-clause, already hidden from OpenCL because C++ for OpenCL 2021 is C++17. MSL 4.1 reports __cplusplus 201703L for the same reason, so it needs the same treatment. Left unhidden the declaration does not parse, and the failure cascades through the rest of the header. With these, SMatrixGPU<float, 5, 5, MatRepSymGPU<float, 5>> compiles as MSL and SMatrixGPU.h itself reports no diagnostics. Host, CUDA, HIP, OpenCL and cling are untouched.
SinCosd and the double specialisation of Abs are the double twins of the float versions, in the same style as the d suffixed helpers in MathUtils; MSL has no double, so they are compiled out there and kept everywhere else. Deterministic mode is refused outright rather than quietly degraded. Its SinCos path computes in double on purpose, for reproducibility against the other backends, and MSL cannot do that at all, so a Metal build that asked for it would silently produce different numbers. It is off by default and opt-in through O2_OVERRIDE_GPUCA_DETERMINISTIC_MODE. The diff adds guards and removes nothing; host, CUDA, HIP, OpenCL and cling are untouched, deterministic mode included.
Three unrelated things, all in the GPU folder. MSL rejects the noexcept specifier. It appears in one header only, GPUCommonAlgorithm.h, so it now goes through GPUnoexcept(), which is noexcept on every other backend and empty on Metal. NDPiecewisePolynomials::getStepWidth and getVertexPosition return double and are not GPUd(), so device code cannot call them in any case; they join the host-only members the file already guards. Spline1DContainerBase::setXrange computes its range width in double for precision. That is device code, so on Metal it uses float instead. This is the one place here where Metal gets a different result rather than simply losing a declaration it could not use. GPUnoexcept() expands to noexcept for host, CUDA, HIP, OpenCL and cling.
Four shapes, all the same underlying point: MSL has no double, and diagnoses it at the declaration, so these break a device translation unit that merely includes them. The double specialisations of sincos, twoPi and pi are guarded; the primary templates still serve float. StatAccumulator has no GPUd() members and is not referenced under GPU/, so the whole struct is host-only. The CircleXYd_t, IntervalXYd_t, Bracketd_t and Rotation2Dd_t aliases are guarded. Bracket.h also used noexcept, which MSL rejects, so it goes through GPUnoexcept() like GPUCommonAlgorithm.h. Nothing is removed for host, CUDA, HIP, OpenCL or cling.
… code The LHC and TPC geometry constants are namespace-scope constexpr double, so on Metal they hit both restrictions at once: no double, and program scope needs an address space. GPUglobalconstexpr() with GPUdoubleValue covers both, and GPUdoubleValue is already double everywhere except Metal, so nothing else moves. LHCBunchSpacingMUS is genuinely read from GPUd() code, which consumes it as float in any case. Propagator's double getFieldXYZ and getBz overloads, PropagatorD, TrackParD and TrackParCovD are all double twins whose float versions remain, so they are compiled out on Metal. TrackUtils computes one local in double inside device code; on Metal that is a float, which together with Spline1DContainerBase::setXrange makes two places where Metal gets a different number rather than simply losing a declaration. With this the Metal translation unit has no double and no noexcept diagnostics left, down from 57. Host, CUDA, HIP, OpenCL and cling keep double throughout.
Without one, GPUDefParametersDefaultsDevice.h falls through to '#error GPU TYPE NOT SET' for GPUCA_GPUTYPE_METAL, which the .metal source defines, so nothing under Definitions/ could be compiled for the backend at all. The column is seeded from OPENCL, which is the other portable backend with no vendor-specific tuning; Metal ends up with the same WARP_SIZE 32 and THREAD_COUNT_DEFAULT 256. Real numbers want measuring on a device once the kernels run. detect_gpu_arch reports METAL alongside the others so the generator emits the block. Every existing architecture comes out byte-identical; the generated device header gains only the Metal block and the architecture comment.
GPUCA_CHOICE routes Metal down the OpenCL arm, and three of those spellings do
not exist in MSL.
nan(uint) is not declared, so QuietNaN uses __builtin_nanf(""), which is what
the CUDA and HIP arm already uses.
remainder() does not exist either; MSL has fmod only. Remainderf is therefore
computed as x - y * rint(x / y), which is the definition of the IEEE remainder
and agrees with remainderf bit for bit over half a million samples across the
range its only caller uses, ITSMFT wrapping an angle difference into TwoPI.
MSL's sincos returns the sine and takes the cosine by thread reference rather
than by pointer, and cannot write through the generic reference SinCos is given,
so the result goes via a local.
Nothing is removed; host, CUDA, HIP, OpenCL and cling keep the GPUCA_CHOICE
arms they had.
Both of these already exist for OpenCL, for reasons that apply unchanged to Metal. The processing settings block in GPUSettingsList.h is skipped for OpenCL because it declares std::string and std::vector members, which GPUSettings.h explicitly does not include for device code. Metal needs the same exclusion. These configs are host-side only: GPUParam carries GPUSettingsRec and GPUSettingsParam, and the processing settings appear only as pointer arguments to host methods, so nothing transferred changes shape. GPUCommonBitSet already carries an extra constructor for OpenCL's __constant. Metal needs the opposite: MSL will not use a user-declared copy constructor to build an object in the constant address space, which is where GPUconstexpr() arrays of bitset live, and leaving the copy constructor implicit makes them constructible again. That one line accounted for 84 of the remaining diagnostics, across DetID and GlobalTrackID. Metal translation unit: 136 errors to 27.
MathUtils kept a using-declaration for StatAccumulator after the struct itself became host-only, which left a dangling name on Metal. GPUCommonAlgorithm::sortOnDevice takes an auto parameter, which is C++20. It is already skipped for OpenCL, at C++17, and MSL 4.1 reports C++17 as well. GPUTPCTrackParam::TransportToXAlpha declares its material constants static at function scope, which MSL rejects; constexpr without static is accepted, as in SMatrixGPU.
The track propagation computes its Jacobian and covariance intermediates in double even when the track is float, because the terms cancel: jj = dx * (dy2dx - f2 * r2inv) is a difference of nearly equal quantities. MSL has no double, and plain float is not an option there. GPUdoubleCalc is compensated two-float arithmetic: the value is mHi + mLo, so the rounding error of each operation is carried explicitly. Over 500k samples of the isolated jj expression, float intermediates reach 3.6e-2 relative error with 23 samples past 1e-3, while the two-float form lands at 5.9e-8, half a float ulp, with none past 1e-3. End to end over 524288 tracks with strongly correlated covariances it keeps the covariance within 1.8e-6 of sqrt(C_ii C_jj) of the CPU double result, where plain float is at 1.2e-5 and the CPU double is itself 1.1e-6 away from a cancellation-free reference. Off Metal it is a plain double, so nothing else moves: substituting the alias back reproduces the previous source exactly, and the files compile unchanged with their real build flags. Cost on an M1 Max, over the whole propagateTo kernel with fast math off: 2.85x the plain float path. Metal has no fp64 at all, so the comparison is against not compiling. The type is defined for every backend so the arithmetic can be tested on the host, where GPUCA_FORCE_DOUBLECALC and GPUCA_FORCE_FLOATCALC select the representation explicitly.
Metal has no ambient work-item builtins: the thread and threadgroup positions exist only as attributes on the kernel entry point, so get_global_id() and its siblings cannot be expressed in a device function the way CUDA's threadIdx or OpenCL's get_global_id() can. Every one of these call sites is inside a function that already receives nBlocks, nThreads, iBlock and iThread, or is one call away from one, so they now use those directly. The substitution is exact on every backend: CUDA, HIP and OpenCL pass precisely these values into Thread(), and the CPU backend passes nThreads = 1 and iThread = 0, which is what the CPU definitions of the macros already assumed. Four helpers had no index in scope and gain one parameter: sortInBlock, GPUTPCCFClusterizer::buildCluster, GPUTPCCFNoiseSuppression::findMinimaAndPeaks and GPUTPCCFPeakFinder::isPeak. GPUCA_SHARED_CACHE and GPUCA_TBB_KERNEL_LOOP likewise take the indices as arguments rather than capturing them from the expansion context. The get_*() macros remain for the kernel entry point in GPUReconstructionKernelMacros.h, which is the one place where Metal does provide them.
Metal defaults to fast math, and fast math reassociates GPUdoubleCalcImpl's compensation terms away: measured on the full propagateTo kernel, the two-float type then costs 1.57x for exactly the accuracy of a plain float. The Metal build therefore follows what the OpenCL one already does: fast math unless GPUCA_DETERMINISTIC_MODE asks for otherwise, and GPUdoubleCalc becomes a plain float when fast math is on. This also takes back the unconditional -fno-fast-math added with the type, which is not free as claimed there: on an ALU-bound kernel it costs the plain float path a factor of two (0.055 -> 0.108 ns/propagate on an M1 Max), and that would have been paid by every Metal kernel, including the clusterer and the sector tracker, which have no double-precision intermediates at all. -fmetal-math-mode=safe would keep the compensation intact for 20% less, but it makes division and square root approximate and the GPU result then stops matching the CPU bit-for-bit (79% of covariance elements instead of 100%), which is the opposite of what deterministic mode is for. -fno-fast-math is also the exact analogue of the CUDA flags used there, --prec-div --prec-sqrt --fmad false.
Three changes, measured on the full propagateTo kernel over 524288 tracks with strongly correlated covariances, against the CPU double result (Apple M1 Max, -fno-fast-math, median of ten runs): * mLo is no longer renormalised after each operation. The value is still mHi + mLo and nothing downstream requires |mLo| <= ulp(mHi)/2, so the results are identical to the last bit: 3.11x -> 2.22x the cost of plain float. * the mixed float forms no longer widen the float operand and run the full two-float operation. Its mLo is zero, but with reassociation disabled the compiler is not allowed to fold those terms away. * float += GPUdoubleCalcImpl now rounds once, as float += double does on the host, rather than converting to float and adding. This is what `c00 += b00 + b00 + a00` resolves to, and it is where the remaining error was: +7% cost, and the covariance error against the CPU double result drops from 1.8e-6 to 5.6e-7 of sqrt(C_ii C_jj), with 86% of elements bit-identical instead of 76%. Net: 3.11x -> 2.37x and closer to the CPU than before. For reference, plain float is at 1.2e-5 on the same metric, and the CPU double result is itself 1.1e-6 away from a cancellation-free reference, so the type is now at the noise floor of the formula it implements. Representations narrower than two floats were measured and are not viable: with four halves, limbs two to four of any value below about one fall into half subnormals, leaving fewer effective bits than a single float (worst error 0.46, 42x the cost), and fixed point Q21.42 reaches 7.4e-6 at 16x.
A reference parameter carries its address space into template deduction, so sincos(T ang, T& s, T& c) called with thread-resident floats deduces T as float from the first argument and as thread float from the other two. OpenCL has the same rule, and the three-parameter overload already there for it works verbatim on Metal.
GPUglobalref() expanded to device, which is what the OpenCL 1 port needed. It no longer matches the code: since the constant address space was dropped, the objects the kernels work on are reached through a generic `this`, so a member pointer or the address of a member is a generic pointer, and MSL will not pass one where a device pointer is wanted. MSL 4.1's generic address space spans device, threadgroup and thread, so leaving the annotation off is correct for every one of these. This mirrors the OpenCL C++ branch, where GPUglobalref() is likewise empty and GPUgeneric() carries the annotation. GPUsharedref() and GPUconstantref() stay explicit: threadgroup is still worth pinning where it is known, and constant is not reachable through a generic pointer at all.
GPUAtomic() expanded to metal::atomic, but the counters it guards are plain integers in structures that the host allocates and transfers, and the code reads them directly outside the atomic operations. That is exactly the CUDA and HIP situation, where GPUAtomic() is a no-op and atomicAdd() takes a plain pointer, so Metal now does the same. MSL's atomic_*_explicit do want metal::atomic, which has the same size and alignment, so the operand is cast at the one point where the operation is issued. The cast has to keep the address space, which a generic pointer does not carry into atomic_*_explicit, hence the two overloads: threadgroup stays threadgroup for the *Shared entry points, anything else resolves to device. GPUTPCCFDecodeZS's shared memory had to be declared GPUsharedref() for that to hold; without it the reference is generic by the time AtomicAddShared sees it.
CAMath::Abs is a template that deduces its parameter rather than taking a float, so a call on a GPUdoubleCalc intermediate selects the primary template, which is declared and never defined. Four call sites in the track propagation do exactly that. Compiling to an object hides it, so the Metal build only fails when the kernels are linked; on the host the type is a double and the question does not arise. Every other CAMath entry point takes a float and is reached through the implicit conversion, so only Abs needs the specialisation.
GPUdoubleCalc now names one of four representations, chosen by GPUCA_DOUBLECALC: hardware double, plain float, the compensated two-float type, or a full IEEE-754 binary64 in software. The default is unchanged, double everywhere except Metal and the two-float type there, so this is a no-op for every existing build. It replaces GPUCA_FORCE_DOUBLECALC and GPUCA_FORCE_FLOATCALC, which could only express two of the four. The binary64 emulation is round to nearest even with subnormals, infinities and NaNs, and correctly rounded division; NaN propagation follows the ARM64 order so an Apple host is a reference down to the payload. It is not for production: over the whole propagateTo kernel on an M1 Max it costs of the order of a hundred times plain float, against roughly two for the two-float type. It is here because it reproduces the CPU result exactly -- every covariance element bit-identical -- which turns a disagreement between Metal and the CPU into a search over code rather than over numerics. Measured on that kernel, 524288 tracks with strongly correlated covariances, median of ten runs, error as |dC_ij| / sqrt(C_ii C_jj) against the CPU double result: plain float 0.108 ns and 1.2e-5, two-float 0.256 ns and 5.6e-7, binary64 12.2 ns and zero. Only the binary64 kernel is unstable run to run, between 9.1 and 12.3 ns; the others vary by a few percent. For scale, the CPU double result is itself 1.1e-6 from a cancellation-free reference, so the two-float type is already at the noise floor of the formula. For scale in the other direction, the same propagateTo on this machine's CPU, where double is native, costs 23.5 ns in float and 27.1 ns in double, so real hardware double is a 1.15x proposition. None of this arithmetic would be needed if Metal had it. Validated against hardware double on 9M operand pairs per operation, over uniform bit patterns, near-equal exponents, the subnormal band, sparse mantissas for the exact and halfway cases, the special values and track-like magnitudes, plus 4M conversions each way: no mismatches. The arithmetic must stay out of line on Metal. Inlined into a kernel it drops the occupancy and runs two to five times slower than the call.
MSL has no static storage duration inside a function. These are all scalar constexpr values used as compile-time constants, so removing static changes nothing for any backend: none of them is odr-used, and no storage was ever emitted for them.
The definitions in GPUTPCGMMergerGPU.cxx qualify smem with GPUsharedref(), and so does every other declaration in the header; this one did not. It makes no difference where GPUsharedref() is empty, but on Metal the declaration then takes a generic reference and the definition a threadgroup one, which are different types.
fragment is a reserved word in MSL, where it qualifies a shader stage, so it cannot name a member, a parameter or a local. The cluster finder used it for all three, which was 35 of the errors left in the Metal compile. Purely a rename, in code positions only: the option descriptions in GPUSettingsList.h that mention fragments are strings and are untouched, as are the comments.
A constructor's implicit `this` is generic in MSL 4.1, and a generic pointer does not reach the constant address space, so a class-type constant at program scope could not be constructed at all: gpustd::bitset for the DetID and GlobalTrackID masks, CfChargePos for INVALID_CHARGE_POS. MSL lets a member function be qualified with the address space of its `this`, so each of these gains a constant-qualified overload next to the existing one, alongside the copy constructor and the OpenCL __constant one that are already there for the same reason. Only the members actually called on a constant object need it.
The shim's partial specialization on a bare T* does match a pointer type written out in full, but not one deduced from an argument, which carries its address space. RDHUtils uses std::is_pointer to keep its pointer overloads apart from its reference ones, so the reference template was instantiated for a pointer and dereferenced it as a struct. metal::is_pointer is address-space aware, so the Metal branch forwards to it.
It forwards straight to PadPlane::getPadRowNumber, which already takes the alias.
MSL supports neither goto nor labels. The label sat at the end of the loop body, so each jump was a continue for the outer loop that could not be written as one because it was issued from an inner loop. The flag is set there instead, breaks out of the loop it was raised in, and continues the outer one; where the jump came from two levels down it breaks twice. The k loop is left early exactly as before, and nothing between the old jumps and the old label ran then either.
The parameter is a reference, so T is deduced with the address space attached, and forcing that same T on a by-value parameter is a substitution failure on Metal. Letting the inner call deduce its own type gives the same T everywhere else, where the address space is not part of it.
The double overloads of getFieldXYZ and getBz were already guarded where they are declared, but not where they are defined. The explicit double in the crossing-point helper becomes GPUdoubleValue, and the differences of nearly equal crossing and centre coordinates that feed atan2 become GPUdoubleCalc, matching the phiCross and dphi next to them. Both aliases are double everywhere except Metal, so the preprocessed source is unchanged for every other backend and for the host.
Unlike ROOT's SMatrix, where the product of two matrices is a matrix, SMatrixGPU returns a lazy expression, and every element of that expression reads the whole of the left operand. Assigning it back element by element therefore reads values that have already been overwritten. It also did not compile: the expression matched the generic operator= that copies the representation, which an expression does not have. This was unreachable until now, since nothing instantiated a matrix multiply assignment in device code.
MatrixDSym5 and MatrixD5 were kept off Metal because SMatrix<double> cannot be named there, which left the track-to-track chi2 and update out of the device build. Spelling them in GPUdoubleCalc brings them back: it is double on every other backend and on the host, so the aliases and everything using them are unchanged there, and on Metal they become the compensated two-float type that the rest of this file already uses. The two remaining explicit doubles go the same way, and the SMatrix written out in full is just MatrixD5.
A kernel's buffers are device memory and the address arrives as an integer, but the Thread() entry points take their pointer arguments unannotated, which is the generic address space. Forwarding a device pointer made the call deduce a device pointer for Args..., which matched no explicit specialisation, so the kernels linked against a Thread() that is declared and never defined. The cast goes through device first rather than straight from the integer, so the generic pointer is formed by the normal conversion.
The Metal frameworks are present on every macOS, so finding them said nothing about whether the backend can be built. It needs MSL 4.1, which arrives with macOS 27 and its toolchain, and a library compiled as MSL 4.1 only loads on macOS 27 and later. The deployment target has no say in this: the -std= flag is what picks the target OS, and MACOSX_DEPLOYMENT_TARGET and -mmacosx-version-min are both ignored by the Metal compiler. Compiling a three-line kernel answers the question directly. AUTO now turns Metal off on an older toolchain instead of failing somewhere in the middle of the build, and an explicit ENABLE_METAL=ON says why it cannot be honoured.
GPUReconstructionMETAL.metal had the device headers and the kernel list behind #if 0, with a comment saying they do not compile as MSL yet. They do now, so the entry point includes them and the library it produces contains the 96 kernels rather than nothing.
Collaborator
davidrohr
reviewed
Sep 21, 2026
| template <> | ||
| GPUdii() void GPUTPCTrackletConstructor::Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& sMem, processorType& GPUrestrict() tracker) | ||
| { | ||
| if (get_local_id(0) == 0) { |
Collaborator
There was a problem hiding this comment.
we support get_global_id and get_local_id. Why don't you just provide them in metal? I think in CUDA we also just provide it via a define.
davidrohr
reviewed
Sep 21, 2026
| if(HIP_ENABLED) | ||
| add_subdirectory(Base/hip) | ||
| endif() | ||
| if(METAL_ENABLED) |
Collaborator
There was a problem hiding this comment.
why do you treat metal differently compared to CUDA and the rest?
davidrohr
reviewed
Sep 21, 2026
| }); \ | ||
| } else { \ | ||
| GPUCA_TBB_KERNEL_LOOP_HOST(rec, vartype, varname, iEnd, code) \ | ||
| #define GPUCA_TBB_KERNEL_LOOP(rec, nBlocks, nThreads, iBlock, iThread, vartype, varname, iEnd, code) \ |
Collaborator
There was a problem hiding this comment.
why do you need to change TBB code for Metal?
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
3 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.
GPUReconstructionMETAL.metal had the device headers and the kernel list behind
#if 0, with a comment saying they do not compile as MSL yet. They do now, so the
entry point includes them and the library it produces contains the 96 kernels
rather than nothing.
Stack created with Sapling. Best reviewed with ReviewStack.