Conversation
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 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.
Same treatment the TPC constants already had. MSL requires every variable at program scope to name an address space, and diagnoses it at the declaration, so a header full of plain constexpr breaks any device translation unit that merely includes it -- whether or not the constants are used. GPUglobalconstexpr() expands to constexpr everywhere except Metal, where it adds constant. Class-scope static members need it too: MSL counts them as program scope. constexpr functions do not, and are left alone. The six vDrift and ExB calibration defaults are double, which does not exist in MSL at all. They are host-only -- nothing under GPU/ refers to them -- so they are now compiled out of device code rather than converted, which would have changed their precision on CUDA and HIP. Preprocessed output is unchanged for host, CUDA, HIP and cling; device code sees the same constants minus those six doubles. Together this takes the TRD headers from 144 errors to 0 in a Metal translation unit.
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.
MSL requires every variable at program scope to name an address space and diagnoses it at the declaration, so these headers broke any device translation unit that included them. GPUglobalconstexpr() expands to constexpr everywhere except Metal, where it adds constant. Preprocessed declarations are unchanged for host, CUDA, HIP and cling.
MSL requires every variable at program scope, class-scope statics included, to name an address space. GPUglobalconstexpr() expands to constexpr everywhere except Metal, so this is a no-op for host, CUDA, HIP, OpenCL and cling.
MSL requires every variable at program scope, class-scope statics included, to name an address space. GPUglobalconstexpr() expands to constexpr everywhere except Metal, so this is a no-op for host, CUDA, HIP, OpenCL and cling.
…ress space MSL requires every variable at program scope, class-scope statics included, to name an address space. GPUglobalconstexpr() expands to constexpr everywhere except Metal, so this is a no-op for host, CUDA, HIP, OpenCL and cling.
MSL requires every variable at program scope, class-scope statics included, to name an address space. GPUglobalconstexpr() expands to constexpr everywhere except Metal, so this is a no-op for host, CUDA, HIP, OpenCL and cling.
…ress space MSL requires every variable at program scope, class-scope statics included, to name an address space. GPUglobalconstexpr() expands to constexpr everywhere except Metal, so this is a no-op for host, CUDA, HIP, OpenCL and cling.
MSL requires every variable at program scope, class-scope statics included, to name an address space. GPUglobalconstexpr() expands to constexpr everywhere except Metal, so this is a no-op for host, CUDA, HIP, OpenCL and cling.
… space MSL requires every variable at program scope, class-scope statics included, to name an address space. GPUglobalconstexpr() expands to constexpr everywhere except Metal, so this is a no-op for host, CUDA, HIP, OpenCL and cling.
MSL requires every variable at program scope, class-scope statics included, to name an address space. GPUglobalconstexpr() expands to constexpr everywhere except Metal, so this is a no-op for host, CUDA, HIP, OpenCL and cling.
MSL requires every variable at program scope, class-scope statics included, to name an address space and diagnoses it at the declaration, so these headers broke any device translation unit that included them regardless of whether the constants were used. GPUglobalconstexpr() expands to constexpr everywhere except Metal, where it adds constant. Last of the series: with this the Metal translation unit has no program-scope diagnostics left, down from 208.
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.
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.
GPUORTFloat16.h is O2's GPU port of the ONNX Runtime float16 types, already carrying 47 GPUd() annotations and already guarding its system includes on GPUCA_GPUCODE_DEVICE, so it is maintained here rather than vendored verbatim. It reaches device code through GPUTPCNNClusterizerKernels.cxx, which converts floats to Float16_t inside GPUd() functions, so it is genuinely device code and not an accidental inclusion. Every one of its 40 diagnostics came from two things. MSL rejects noexcept, and because the failure is mid-declaration the constructor bodies then lost their member names, which is where the 'undeclared identifier val', the 'protected member' complaints and the sizeof static_assert came from. The rest were class-scope constexpr needing the constant address space. Metal translation unit: 1157 errors to 1108.
Two things MSL does differently from every other backend. It has no work-item builtins: the grid dimensions arrive as kernel attributes. Without a branch of its own Metal fell through to the host definitions of get_group_id() and friends, which name iBlock and nBlocks, variables that only exist in the host-side loop. And it requires every kernel parameter to carry an attribute, so the sector index cannot be passed by value. GPUCA_KRNLGPU_DEF therefore gets two hooks, GPUCA_KRNL_SECTOR_ARG and GPUCA_KRNL_GRID_ARGS, which the Metal source fills in with a buffer and the four grid attributes. Both default to what the signature had, so CUDA, HIP and OpenCL generate exactly the same entry point as before. Kernel list diagnostics: 408 to 96, and the translation unit 1108 to 998. The 96 left are the 48 kernels that take arguments, which still need an answer for how Metal passes them.
MSL needs an explicit, distinct buffer index on every kernel parameter, and the preprocessor cannot supply one: the kernel list splices arguments as a flat comma-separated list, and __COUNTER__ is monotonic across the translation unit rather than per kernel. o2_gpu_add_kernel already walks the arguments in pairs, so it emits the index there instead. Indices start at 3, after gpu_mem, the constant memory and the sector. Declarations now go through GPUPtr1(idx, type, name) for pointers and GPUArg1(idx, type, name) for scalars, which each backend defines as it needs. Metal masks pointers as a 64-bit address exactly as OpenCL does, and for the same reason: GPUTRDTrackerKernels takes a GPUTRDTrackerGPU*, and a pointer to a derived class is not a valid kernel argument type there either. Binding POD pointers directly would have worked but would not have covered that case, so both go the same way. Generated entry points are byte-identical for CUDA, HIP and OpenCL. Kernel list diagnostics: 408 to 0, and the translation unit 1108 to 881.
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.
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.
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.
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... |
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.