1// Metal backend host code (Objective-C++). See shaders/raid_kernels.metal for the kernels and the precision design.2#import <Metal/Metal.h>3#import <Foundation/Foundation.h>4#include "raid/backend/backend.hpp"5#include "raid/core/error.hpp"6#include "raid/core/log.hpp"7#include "raid/statistics/reductions.hpp"8#include <chrono>9#include <cmath>10#include <map>11#include <mutex>12#include <unistd.h>13#include <mach-o/dyld.h>1415namespace raid { namespace metal { const char* msl_bundle(); } }1617namespace raid::backend {1819namespace {20struct Chunk { uint32_t seg, begin, end, pad; };21struct Timer { std::chrono::steady_clock::time_point t0 = std::chrono::steady_clock::now(); f64 s() const { return std::chrono::duration<f64>(std::chrono::steady_clock::now() - t0).count(); } };22}2324class MetalBackend final : public Backend {25 public:26 explicit MetalBackend(PrecisionPolicy policy) : policy_(policy) {27 @autoreleasepool {28 device_ = MTLCreateSystemDefaultDevice();29 if (!device_) { log::warn("metal: no device"); return; }30 queue_ = [device_ newCommandQueue];31 available_ = device_ != nil && queue_ != nil;32 }33 }34 ~MetalBackend() override = default;35 std::string name() const override { return "metal"; }36 raid::Backend kind() const override { return raid::Backend::metal; }37 bool available() const override { return available_; }38 Device device() const override {39 Device d; d.kind = raid::Backend::metal; d.supports_fp64 = false; d.unified_memory = device_.hasUnifiedMemory;40 d.name = device_ ? std::string([device_.name UTF8String]) : "none"; d.memory_bytes = device_.recommendedMaxWorkingSetSize; d.max_buffer_bytes = device_.maxBufferLength;41 d.compute_units = 0; // Metal does not expose the GPU core count; recorded from system_profiler in the manifest42 d.details = std::string("families:") + ([device_ supportsFamily:MTLGPUFamilyApple9] ? "Apple9" : ([device_ supportsFamily:MTLGPUFamilyApple8] ? "Apple8" : "Apple7-")) + " metal3:" + ([device_ supportsFamily:MTLGPUFamilyMetal3] ? "yes" : "no") + " library:" + library_source_;43 return d;44 }4546 // ---- library / pipelines ----------------------------------------------------------------------------------------47 bool ensure_library(KernelStats* st) {48 std::lock_guard<std::mutex> lk(m_);49 if (library_) return true;50 @autoreleasepool {51 NSError* err = nil; Timer t;52 // 1. prebuilt metallib: RAID_METALLIB env, compile-time path, next to the executable53 std::vector<std::string> candidates;54 if (const char* e = std::getenv("RAID_METALLIB")) candidates.push_back(e);55#ifdef RAID_METALLIB_PATH56 candidates.push_back(RAID_METALLIB_PATH);57#endif58 { char buf[4096]; uint32_t sz = sizeof buf; if (_NSGetExecutablePath(buf, &sz) == 0) { std::string p(buf); auto k = p.rfind('/'); if (k != std::string::npos) candidates.push_back(p.substr(0, k) + "/raid.metallib"); } }59 for (auto& c : candidates) {60 if (access(c.c_str(), R_OK) != 0) continue;61 library_ = [device_ newLibraryWithURL:[NSURL fileURLWithPath:[NSString stringWithUTF8String:c.c_str()]] error:&err];62 if (library_) { library_source_ = "metallib:" + c; compile_seconds_ = t.s(); break; }63 }64 // 2. runtime compilation of the embedded MSL65 if (!library_) {66 MTLCompileOptions* opts = [MTLCompileOptions new]; opts.fastMathEnabled = NO; // keep IEEE semantics (no reassociation): required by the double-float transforms67 library_ = [device_ newLibraryWithSource:[NSString stringWithUTF8String:raid::metal::msl_bundle()] options:opts error:&err];68 if (!library_) { log::error("metal: shader compilation failed", {{"error", err ? std::string([[err localizedDescription] UTF8String]) : "?"}}); available_ = false; return false; }69 library_source_ = "runtime"; compile_seconds_ = t.s(); if (st) st->compiled_now = true;70 }71 }72 return true;73 }74 id<MTLComputePipelineState> pipeline(const std::string& fn) {75 std::lock_guard<std::mutex> lk(m_);76 auto it = pipelines_.find(fn); if (it != pipelines_.end()) return it->second;77 @autoreleasepool {78 NSError* err = nil; id<MTLFunction> f = [library_ newFunctionWithName:[NSString stringWithUTF8String:fn.c_str()]];79 RAID_REQUIRE(f != nil, unsupported, "metal: kernel not found: " + fn);80 id<MTLComputePipelineState> p = [device_ newComputePipelineStateWithFunction:f error:&err];81 RAID_REQUIRE(p != nil, unsupported, "metal: pipeline failed for " + fn + ": " + (err ? std::string([[err localizedDescription] UTF8String]) : "?"));82 pipelines_[fn] = p; return p;83 }84 }85 // ---- buffer pool (shared storage; unified memory) ---------------------------------------------------------------86 id<MTLBuffer> buffer(usize bytes, usize* allocs) {87 std::lock_guard<std::mutex> lk(m_);88 bytes = std::max<usize>(bytes, 16); usize cls = 1; while (cls < bytes) cls <<= 1;89 auto& v = pool_[cls]; if (!v.empty()) { auto b = v.back(); v.pop_back(); return b; }90 if (allocs) ++*allocs; ++total_allocs_;91 id<MTLBuffer> b = [device_ newBufferWithLength:cls options:MTLResourceStorageModeShared];92 RAID_REQUIRE(b != nil, memory, "metal: buffer allocation failed (" + std::to_string(cls) + " bytes)");93 return b;94 }95 void release(id<MTLBuffer> b) { std::lock_guard<std::mutex> lk(m_); pool_[b.length].push_back(b); }96 usize total_allocations() const { return total_allocs_; }9798 // ---- helpers -----------------------------------------------------------------------------------------------------99 std::vector<Chunk> make_chunks(std::span<const i64> offsets, usize chunk_max, usize* n_segments) const {100 std::vector<Chunk> ch; const usize S = offsets.size() - 1; if (n_segments) *n_segments = S;101 for (usize s = 0; s < S; ++s) { const u64 b = static_cast<u64>(offsets[s]), e = static_cast<u64>(offsets[s + 1]); RAID_REQUIRE(e <= 0xFFFFFFFFull, unsupported, "metal: segments beyond 2^32 elements"); if (e == b) { ch.push_back({static_cast<uint32_t>(s), static_cast<uint32_t>(b), static_cast<uint32_t>(e), 0}); continue; }102 for (u64 x = b; x < e; x += chunk_max) ch.push_back({static_cast<uint32_t>(s), static_cast<uint32_t>(x), static_cast<uint32_t>(std::min<u64>(e, x + chunk_max)), 0}); }103 return ch;104 }105 usize chunk_size() const { return policy_.fp32_chunk_max == 0 ? 4096 : policy_.fp32_chunk_max; }106 static void finish(KernelStats* st, const char* k, const Timer& t, id<MTLCommandBuffer> cb, usize bin, usize bout, usize allocs, const std::string& note) {107 if (!st) return; st->kernel = k; st->backend = raid::Backend::metal; st->precision = Precision::fp32; st->wall_seconds = t.s(); st->device_seconds = cb ? (cb.GPUEndTime - cb.GPUStartTime) : 0; st->bytes_in = bin; st->bytes_out = bout; st->buffer_allocations = allocs; st->note = note;108 }109 // Fused realized kernel; returns 5 vectors combined in FP64 on the CPU (adds the chunk-boundary |x_i||x_{i-1}| terms).110 void realized_all(const Batch& b, std::vector<f64>& s2, std::vector<f64>& adj, std::vector<f64>& neg, std::vector<f64>& pos, std::vector<f64>& s4, KernelStats* st, int single_op = -1) {111 Timer t; usize allocs = 0; RAID_REQUIRE(available_ && ensure_library(st), unsupported, "metal unavailable");112 usize S; auto chunks = make_chunks(b.offsets, chunk_size(), &S); const usize n = b.values.size(), C = chunks.size();113 s2.assign(S, 0); adj.assign(S, 0); neg.assign(S, 0); pos.assign(S, 0); s4.assign(S, 0);114 if (n == 0 || C == 0) { finish(st, "realized_fused", t, nil, 0, 0, 0, library_source_); return; }115 @autoreleasepool {116 id<MTLBuffer> bx = buffer(n * 4, &allocs), bc = buffer(C * sizeof(Chunk), &allocs), bo = buffer(C * 5 * 8, &allocs);117 Timer tconv; float* xf = static_cast<float*>(bx.contents); for (usize i = 0; i < n; ++i) xf[i] = static_cast<float>(b.values[i]); const f64 conv_s = tconv.s();118 std::memcpy(bc.contents, chunks.data(), C * sizeof(Chunk));119 id<MTLCommandBuffer> cb = [queue_ commandBuffer]; id<MTLComputeCommandEncoder> enc = [cb computeCommandEncoder];120 const uint32_t op = single_op < 0 ? 0 : static_cast<uint32_t>(single_op);121 id<MTLComputePipelineState> p = pipeline(single_op < 0 ? "realized_fused" : "segment_reduce");122 [enc setComputePipelineState:p]; [enc setBuffer:bx offset:0 atIndex:0]; [enc setBuffer:bc offset:0 atIndex:1]; [enc setBuffer:bo offset:0 atIndex:2]; if (single_op >= 0) [enc setBytes:&op length:4 atIndex:3];123 [enc dispatchThreadgroups:MTLSizeMake(C, 1, 1) threadsPerThreadgroup:MTLSizeMake(256, 1, 1)]; [enc endEncoding]; [cb commit]; [cb waitUntilCompleted];124 RAID_REQUIRE(cb.status == MTLCommandBufferStatusCompleted, numerical, "metal: command buffer failed");125 const float* o = static_cast<const float*>(bo.contents);126 for (usize c = 0; c < C; ++c) { const usize s = chunks[c].seg;127 if (single_op < 0) { s2[s] += static_cast<f64>(o[c * 10 + 0]) + o[c * 10 + 1]; adj[s] += static_cast<f64>(o[c * 10 + 2]) + o[c * 10 + 3]; neg[s] += static_cast<f64>(o[c * 10 + 4]) + o[c * 10 + 5]; pos[s] += static_cast<f64>(o[c * 10 + 6]) + o[c * 10 + 7]; s4[s] += static_cast<f64>(o[c * 10 + 8]) + o[c * 10 + 9]; }128 else { f64 v = static_cast<f64>(o[c * 2]) + o[c * 2 + 1]; (single_op == 0 ? s2 : single_op == 1 ? adj : single_op == 2 ? neg : single_op == 3 ? pos : s4)[s] += v; }129 // chunk boundary term for the adjacent product (FP64 on the CPU)130 if ((single_op < 0 || single_op == 1) && chunks[c].begin > static_cast<uint32_t>(b.offsets[s])) adj[s] += std::abs(b.values[chunks[c].begin]) * std::abs(b.values[chunks[c].begin - 1]); }131 finish(st, single_op < 0 ? "realized_fused" : "segment_reduce", t, cb, n * 8, S * 40, allocs, library_source_ + " chunks=" + std::to_string(C) + " convert_s=" + std::to_string(conv_s));132 release(bx); release(bc); release(bo);133 }134 }135136 // ---- Backend interface ---------------------------------------------------------------------------------------------137 std::vector<f64> sum_squares(const Batch& b, KernelStats* st) override { std::vector<f64> a, c, d, e, f; realized_all(b, a, c, d, e, f, st, 0); return a; }138 std::vector<f64> sum_abs_adjacent(const Batch& b, KernelStats* st) override { std::vector<f64> a, c, d, e, f; realized_all(b, a, c, d, e, f, st, 1); return c; }139 void semivariances(const Batch& b, std::vector<f64>& neg, std::vector<f64>& pos, KernelStats* st) override { std::vector<f64> a, c, f; realized_all(b, a, c, neg, pos, f, st, -1); }140 std::vector<f64> sum_fourth(const Batch& b, KernelStats* st) override { std::vector<f64> a, c, d, e, f; realized_all(b, a, c, d, e, f, st, 4); return f; }141 /// All five realized accumulators in a single fused pass (used by est::realized_measures_batch when the backend is Metal).142 void realized_fused(const Batch& b, std::vector<f64>& s2, std::vector<f64>& adj, std::vector<f64>& neg, std::vector<f64>& pos, std::vector<f64>& s4, KernelStats* st) { realized_all(b, s2, adj, neg, pos, s4, st, -1); }143144 std::vector<f64> rolling_sum_squares(const Batch& b, usize window, KernelStats* st) override {145 Timer t; usize allocs = 0; RAID_REQUIRE(available_ && ensure_library(st), unsupported, "metal unavailable");146 const usize n = b.values.size(), S = b.n_series(); std::vector<f64> out(n, std::numeric_limits<f64>::quiet_NaN()); if (n == 0) return out;147 RAID_REQUIRE(window <= 65536, unsupported, "metal rolling window too large for the direct kernel");148 @autoreleasepool {149 id<MTLBuffer> bx = buffer(n * 4, &allocs), bb = buffer(S * 4, &allocs), bs = buffer(n * 4, &allocs), bo = buffer(n * 4, &allocs);150 float* xf = static_cast<float*>(bx.contents); for (usize i = 0; i < n; ++i) xf[i] = static_cast<float>(b.values[i]);151 uint32_t* sb = static_cast<uint32_t*>(bb.contents); uint32_t* so = static_cast<uint32_t*>(bs.contents);152 for (usize s = 0; s < S; ++s) { sb[s] = static_cast<uint32_t>(b.offsets[s]); for (i64 i = b.offsets[s]; i < b.offsets[s + 1]; ++i) so[i] = static_cast<uint32_t>(s); }153 const uint32_t w = static_cast<uint32_t>(window), nn = static_cast<uint32_t>(n);154 id<MTLCommandBuffer> cb = [queue_ commandBuffer]; id<MTLComputeCommandEncoder> enc = [cb computeCommandEncoder]; auto p = pipeline("rolling_sumsq");155 [enc setComputePipelineState:p]; [enc setBuffer:bx offset:0 atIndex:0]; [enc setBuffer:bb offset:0 atIndex:1]; [enc setBuffer:bs offset:0 atIndex:2]; [enc setBuffer:bo offset:0 atIndex:3]; [enc setBytes:&w length:4 atIndex:4]; [enc setBytes:&nn length:4 atIndex:5];156 [enc dispatchThreads:MTLSizeMake(n, 1, 1) threadsPerThreadgroup:MTLSizeMake(std::min<usize>(256, p.maxTotalThreadsPerThreadgroup), 1, 1)]; [enc endEncoding]; [cb commit]; [cb waitUntilCompleted];157 RAID_REQUIRE(cb.status == MTLCommandBufferStatusCompleted, numerical, "metal: command buffer failed");158 const float* o = static_cast<const float*>(bo.contents); for (usize i = 0; i < n; ++i) out[i] = o[i];159 finish(st, "rolling_sumsq", t, cb, n * 8, n * 8, allocs, library_source_); release(bx); release(bb); release(bs); release(bo);160 }161 return out;162 }163164 void regression_stats(std::span<const f64> x, std::span<const f64> y, std::span<const i64> offsets, usize k, std::vector<f64>& xtx, std::vector<f64>& xty, std::vector<f64>& yty, KernelStats* st) override {165 Timer t; usize allocs = 0; RAID_REQUIRE(available_ && ensure_library(st), unsupported, "metal unavailable");166 RAID_REQUIRE(k >= 1 && k <= 8, unsupported, "metal regression_stats supports k <= 8");167 usize S; auto chunks = make_chunks(offsets, chunk_size(), &S); const usize rows = y.size(), C = chunks.size(); const usize ntri = k * (k + 1) / 2, nstat = ntri + k + 1;168 xtx.assign(S * k * k, 0); xty.assign(S * k, 0); yty.assign(S, 0); if (rows == 0) return;169 @autoreleasepool {170 id<MTLBuffer> bX = buffer(rows * k * 4, &allocs), bY = buffer(rows * 4, &allocs), bc = buffer(C * sizeof(Chunk), &allocs), bo = buffer(C * nstat * 8, &allocs);171 float* xf = static_cast<float*>(bX.contents); for (usize i = 0; i < rows * k; ++i) xf[i] = static_cast<float>(x[i]); float* yf = static_cast<float*>(bY.contents); for (usize i = 0; i < rows; ++i) yf[i] = static_cast<float>(y[i]);172 std::memcpy(bc.contents, chunks.data(), C * sizeof(Chunk)); const uint32_t kk = static_cast<uint32_t>(k);173 id<MTLCommandBuffer> cb = [queue_ commandBuffer]; id<MTLComputeCommandEncoder> enc = [cb computeCommandEncoder]; auto p = pipeline("regression_stats");174 [enc setComputePipelineState:p]; [enc setBuffer:bX offset:0 atIndex:0]; [enc setBuffer:bY offset:0 atIndex:1]; [enc setBuffer:bc offset:0 atIndex:2]; [enc setBuffer:bo offset:0 atIndex:3]; [enc setBytes:&kk length:4 atIndex:4];175 [enc dispatchThreadgroups:MTLSizeMake(C, 1, 1) threadsPerThreadgroup:MTLSizeMake(256, 1, 1)]; [enc endEncoding]; [cb commit]; [cb waitUntilCompleted];176 RAID_REQUIRE(cb.status == MTLCommandBufferStatusCompleted, numerical, "metal: command buffer failed");177 const float* o = static_cast<const float*>(bo.contents);178 for (usize c = 0; c < C; ++c) { const usize s = chunks[c].seg; usize q = 0;179 for (usize a = 0; a < k; ++a) for (usize bcol = a; bcol < k; ++bcol) { const f64 v = static_cast<f64>(o[(c * nstat + q) * 2]) + o[(c * nstat + q) * 2 + 1]; xtx[s * k * k + a * k + bcol] += v; if (a != bcol) xtx[s * k * k + bcol * k + a] += v; ++q; }180 for (usize a = 0; a < k; ++a) { xty[s * k + a] += static_cast<f64>(o[(c * nstat + q) * 2]) + o[(c * nstat + q) * 2 + 1]; ++q; }181 yty[s] += static_cast<f64>(o[(c * nstat + q) * 2]) + o[(c * nstat + q) * 2 + 1]; }182 finish(st, "regression_stats", t, cb, rows * (k + 1) * 8, S * (k * k + k + 1) * 8, allocs, library_source_ + " chunks=" + std::to_string(C)); release(bX); release(bY); release(bc); release(bo);183 }184 }185186 std::vector<f64> garch_nll_batch(const Batch& r, std::span<const f64> params, usize n_params, usize p, int model_code, f64 nu_fixed, KernelStats* st) override {187 Timer t; usize allocs = 0; RAID_REQUIRE(available_ && ensure_library(st), unsupported, "metal unavailable");188 const usize S = r.n_series(), n = r.values.size(), J = S * n_params; std::vector<f64> out(J, std::numeric_limits<f64>::infinity()); if (J == 0) return out;189 RAID_REQUIRE(p == 4 || p == 5 || p == 6, invalid_argument, "garch params per vector must be 4 (garch), 5 (gjr/egarch) or 6 (with nu)");190 @autoreleasepool {191 id<MTLBuffer> br = buffer(n * 4, &allocs), bb = buffer(S * 4, &allocs), be = buffer(S * 4, &allocs), bp = buffer(J * 8 * 4, &allocs), bv = buffer(S * 4, &allocs), bj = buffer(J * 16, &allocs), bo = buffer(J * 4, &allocs);192 float* rf = static_cast<float*>(br.contents); for (usize i = 0; i < n; ++i) rf[i] = static_cast<float>(r.values[i]);193 uint32_t* sb = static_cast<uint32_t*>(bb.contents); uint32_t* se = static_cast<uint32_t*>(be.contents); float* v0 = static_cast<float*>(bv.contents);194 float* pf = static_cast<float*>(bp.contents); uint32_t* jb = static_cast<uint32_t*>(bj.contents);195 const bool gjr = model_code == 1, eg = model_code == 2;196 for (usize s = 0; s < S; ++s) { sb[s] = static_cast<uint32_t>(r.offsets[s]); se[s] = static_cast<uint32_t>(r.offsets[s + 1]);197 for (usize q = 0; q < n_params; ++q) { const f64* th = params.data() + (s * n_params + q) * p; const usize j = s * n_params + q;198 // layout [mu, omega, alpha, gamma, beta, nu]; FP64 initial variance per (series, mu) computed on the CPU → we use mu of the first param set per series199 pf[j * 8 + 0] = static_cast<float>(th[0]); pf[j * 8 + 1] = static_cast<float>(th[1]); pf[j * 8 + 2] = static_cast<float>(th[2]);200 pf[j * 8 + 3] = (gjr || eg) ? static_cast<float>(th[3]) : 0.0f; pf[j * 8 + 4] = static_cast<float>(th[(gjr || eg) ? 4 : 3]);201 const bool has_nu = p == 6 || (p == 5 && !gjr && !eg); const f64 nu = has_nu ? th[p - 1] : (nu_fixed > 2 ? nu_fixed : 0.0); pf[j * 8 + 5] = static_cast<float>(nu);202 // FP64 density constants (Metal has no lgamma): kconst and E|z|203 const f64 pi = 3.14159265358979323846;204 if (nu > 2.0) { const f64 lg = std::lgamma((nu + 1) / 2) - std::lgamma(nu / 2); pf[j * 8 + 6] = static_cast<float>(lg - 0.5 * std::log(pi * (nu - 2.0))); pf[j * 8 + 7] = static_cast<float>(2.0 * std::sqrt(nu - 2.0) * std::exp(lg) / ((nu - 1.0) * std::sqrt(pi))); }205 else { pf[j * 8 + 6] = static_cast<float>(-0.5 * std::log(2.0 * pi)); pf[j * 8 + 7] = static_cast<float>(std::sqrt(2.0 / pi)); }206 jb[j * 4 + 0] = static_cast<uint32_t>(s); jb[j * 4 + 1] = static_cast<uint32_t>(j); jb[j * 4 + 2] = 0; jb[j * 4 + 3] = 0; }207 // initial variance: sample variance of (r − mu) with mu from the first parameter vector of the series (FP64)208 const f64 mu = params[(s * n_params) * p]; auto x = r.series(s); f64 m1 = 0; for (f64 v : x) m1 += v - mu; m1 /= static_cast<f64>(x.size()); f64 var0 = 0; for (f64 v : x) var0 += (v - mu - m1) * (v - mu - m1); var0 /= static_cast<f64>(x.size()); v0[s] = static_cast<float>(var0); }209 const uint32_t model = static_cast<uint32_t>(model_code), nj = static_cast<uint32_t>(J);210 id<MTLCommandBuffer> cb = [queue_ commandBuffer]; id<MTLComputeCommandEncoder> enc = [cb computeCommandEncoder]; auto pl = pipeline("garch_nll");211 [enc setComputePipelineState:pl]; [enc setBuffer:br offset:0 atIndex:0]; [enc setBuffer:bb offset:0 atIndex:1]; [enc setBuffer:be offset:0 atIndex:2]; [enc setBuffer:bp offset:0 atIndex:3]; [enc setBuffer:bv offset:0 atIndex:4]; [enc setBuffer:bj offset:0 atIndex:5]; [enc setBuffer:bo offset:0 atIndex:6]; [enc setBytes:&model length:4 atIndex:7]; [enc setBytes:&nj length:4 atIndex:8];212 [enc dispatchThreads:MTLSizeMake(J, 1, 1) threadsPerThreadgroup:MTLSizeMake(std::min<usize>(64, pl.maxTotalThreadsPerThreadgroup), 1, 1)]; [enc endEncoding]; [cb commit]; [cb waitUntilCompleted];213 RAID_REQUIRE(cb.status == MTLCommandBufferStatusCompleted, numerical, "metal: command buffer failed");214 const float* o = static_cast<const float*>(bo.contents); for (usize j = 0; j < J; ++j) out[j] = o[j];215 finish(st, "garch_nll", t, cb, n * 8 + J * p * 8, J * 8, allocs, library_source_ + " jobs=" + std::to_string(J)); release(br); release(bb); release(be); release(bp); release(bv); release(bj); release(bo);216 }217 return out;218 }219 f64 compile_seconds() const { return compile_seconds_; }220 std::string library_source() const { return library_source_; }221222 private:223 PrecisionPolicy policy_; id<MTLDevice> device_ = nil; id<MTLCommandQueue> queue_ = nil; id<MTLLibrary> library_ = nil; bool available_{false};224 std::map<std::string, id<MTLComputePipelineState>> pipelines_; std::map<usize, std::vector<id<MTLBuffer>>> pool_; std::mutex m_; usize total_allocs_{0};225 std::string library_source_{"none"}; f64 compile_seconds_{0};226};227228std::shared_ptr<Backend> make_metal_backend(const PrecisionPolicy& policy) {229 static std::mutex m; static std::weak_ptr<MetalBackend> cached; std::lock_guard<std::mutex> lk(m);230 if (auto p = cached.lock()) return p; // one device/queue/library per process (pipelines and buffers are reused)231 auto p = std::make_shared<MetalBackend>(policy); cached = p; if (!p->available()) return nullptr; return p;232}233234/// Fused realized accumulators when the backend is Metal (falls back to the four generic calls otherwise).235bool realized_fused(Backend& be, const Batch& b, std::vector<f64>& s2, std::vector<f64>& adj, std::vector<f64>& neg, std::vector<f64>& pos, std::vector<f64>& s4, KernelStats* st) {236 if (auto* m = dynamic_cast<MetalBackend*>(&be)) { m->realized_fused(b, s2, adj, neg, pos, s4, st); return true; }237 return false;238}239f64 metal_compile_seconds(Backend& be) { if (auto* m = dynamic_cast<MetalBackend*>(&be)) return m->compile_seconds(); return 0; }240241} // namespace raid::backend242