#2937 Metal build broken on current main (v0.1.45-metal): 3 separate defects in the perf-merge sync

Descrição

# Metal build broken on current main (v0.1.45-metal): 3 separate defects in the perf-merge sync

## Summary

A fresh install on an Apple Silicon Mac (`git clone` + `./setup.sh`, following `docs/MAC.md`) fails to build the engine at commit `477ac13c`. Three separate defects, all introduced because the f-145 perf-PR group merges (`5a6f4ed5`, `8bdef963`, `df6262f2`, `21c57972`, `b277412e`, `cbc6e3d8`, `e8b1edbc`, `f260925a`) and the #2666 peer-experts work changed the shared host code and the kernel headers without re-syncing the Metal side:

1. The Metal compat shim (`metal/include/metal_compat/cuda_runtime.h`) lacks six names the shared host code now uses → **compile errors**.
2. The generated stub file (`metal/src/engine/stubs.cpp`) is stale: one stale signature (member → hard compile error) and 21 entry points that are referenced but undefined → **link failure**.
3. One of the 21, `strata::kernels::fused_gr_persist_init`, is called **unconditionally** from `Verifier::init` (src/core/verify.cpp:646), so even with the stub the engine would stop at start-up (exit 86, "no Metal implementation yet").

The Metal target (`metal/CMakeLists.txt`) is a separate CMake project that never compiles in the CUDA/HIP/SYCL builds, so none of the PC-side changes were caught: the perf group is measured on RTX 2080 Ti / R9700 / gfx12 hardware only.

## Environment

- Apple Silicon Mac, macOS 27.0, Command Line Tools (Apple clang)
- `git clone https://github.com/Niko1221/Strata`, commit `477ac13c` (v0.1.45-metal)
- `./setup.sh` → step 4 (engine build) fails; the model download never starts

## Defect 1 — shim missing six names

`src/core/peer_experts.cpp` (added by #2666; `e6ad6c81` fixed the *HIP* compat for exactly this, "peer_experts.cpp did not compile on HIP in f-145-int", but not the Metal one):

```
peer_experts.cpp:72:13:  error: use of undeclared identifier 'cudaMemcpyPeer'
peer_experts.cpp:157:29: error: use of undeclared identifier 'cudaDeviceDisablePeerAccess'
peer_experts.cpp:158:34: error: use of undeclared identifier 'cudaDeviceDisablePeerAccess'
```

`src/core/verify.cpp` (sm75/cuda86 group):

```
verify.cpp:285:54: error: use of undeclared identifier 'cudaDevAttrComputeCapabilityMajor'
verify.cpp:286:54: error: use of undeclared identifier 'cudaDevAttrComputeCapabilityMinor'
verify.cpp:859:17: error: use of undeclared identifier 'cudaDeviceGetStreamPriorityRange'
verify.cpp:860:17: error: use of undeclared identifier 'cudaStreamCreateWithPriority'
```

(`cudaDeviceDisablePeerAccess` is called with one argument at verify/peer_experts sites; the real CUDA signature has a defaulted `flags`, so the shim needs the default too.)

## Defect 2 — stale generated stubs

After the shim is fixed, compilation passes and the link fails with 21 undefined symbols (ld64):

```
strata::kernels::wait_handoff                     strata::kernels::hc_q8_to_bf16
strata::kernels::handoff_publish                  strata::kernels::gdn_step_norm_groups
strata::kernels::wait_flag_ge_stamped             strata::kernels::fused_gr_persist_init
strata::kernels::fused_gr_read_persist            strata::kernels::qsa_block_scores_rows
strata::kernels::qsa_decode_attn_split            strata::kernels::native_quantize_q8_1_il
strata::kernels::wait_copy_i32_from_mapped        strata::kernels::qsa_decode_attn_batch_rows
strata::kernels::copy_rows_from_mapped_after      strata::kernels::cpu::iq256_gu_rows_exact_one
strata::kernels::res_patch                        strata::prefill::f16_range_read
strata::prefill::gemm_f16_range_read              strata::prefill::gr_norm
strata::prefill::gr_norm_rs                       strata::prefill::gr_write_norm_rs
strata::prefill::gr_write_cvec_norm_rs            strata::prefill::Gemm::bf16 (stale signature)
```

All referenced from the shared host objects (`verify.cpp.o`, `prefill.cpp.o`, `mtp.cpp.o`, `native_expert.cpp.o`, `generate.cpp.o`) — the same calls run on every backend, and the last stub regeneration predates the perf group, so none of them have a Metal implementation or a stub.

The stale-signature part: `Gemm::bf16` gained a trailing `bool x_f16` (`8bdef963`, sm75 series), so the committed 9-parameter out-of-line stub no longer matches the declaration:

```
stubs.cpp:328:29: error: out-of-line definition of 'bf16' does not match any declaration in 'strata::prefill::Gemm'
```

The same drift, in its quiet form, hit four free functions: the committed stubs of `gr_norm` / `gr_norm_rs` / `gr_write_norm_rs` / `gr_write_cvec_norm_rs` are one parameter short of the current `prefill/kernels.hpp` declarations. Free functions make that *compile* (it becomes a dead overload), but the 9/9/12/16-parameter versions the host code now references are what the link needs — and the dead overloads are a latent landmine (an 8-argument call would silently bind to the stub).

## Defect 3 — `fused_gr_persist_init` stops the engine at start

`Verifier::init` calls `strata::kernels::fused_gr_persist_init();` unconditionally (src/core/verify.cpp:646, no `#if`). The feature is `STRATA_HC_PERSIST=1`, a HIP-only read variant: the non-HIP CUDA build compiles an empty body for it (`src/kernels/cuda/fused_gr.cu:3030`, `#else` branch), and on Metal the .cu is not built at all. So on Metal the call reaches the stub and `missing()` stops the engine (metal/src/engine/missing.cpp: "no Metal implementation yet", exit 86) on *every* start, regardless of model or settings. The Metal side needs the same empty body the CUDA build has (the project already does this for the sibling `fused_gr_check()` at metal/src/engine/setup_host.cpp:59).

## Runtime reachability of the 21 (so the stubs never fire on a default Mac)

Audited each call site: 20 of the 21 sit behind opt-in or PC-only gates and are not reached on a default Apple Silicon run —

| entry point | gate |
|---|---|
| `fused_gr_read_persist` | branch requires `fused_gr_supported()`, which on Metal is the stub → `false` |
| `native_quantize_q8_1_il` | `qil` = `STRATA_QIL` env or `device_is_sm75()` (Turing; false on Apple silicon) |
| `gdn_step_norm_groups` (×2) | `STRATA_BATCH_GDN_GROUPS=1` (opt-in) |
| `qsa_block_scores_rows`, `qsa_decode_attn_batch_rows` | `STRATA_BATCH_ATTN_ROWS=1` (opt-in) |
| `qsa_decode_attn_split` | `STRATA_QSA_SPLIT_ATTN=1` (opt-in) |
| `wait_copy_i32_from_mapped`, `copy_rows_from_mapped_after` | `fuse_wait_on()`: `STRATA_FUSE_WAIT` on non-HIP, off by default |
| `wait_flag_ge_stamped` (×3) | `m_bal_ != nullptr` (`set_pcie_balance`, multi-GPU) or the fused-wait branch |
| `wait_handoff`, `handoff_publish` | `hflag_*_d_ != nullptr` (overlapped layer split, multi-GPU) |
| `res_patch` | `drive.d.res_dirty`, set only with `STRATA_ADAPT_EVICT_SYNC` |
| `iq256_gu_rows_exact_one` | `STRATA_CPU_IQ_ALL_EXACT1` **and** `ggml_cpu_has_avx2()` (never true on arm64) |
| `gr_norm`, `gr_norm_rs`, `gr_write_norm_rs`, `gr_write_cvec_norm_rs`, `f16_range_read`, `gemm_f16_range_read`, `hc_q8_to_bf16` | all called only from `Prefill::run_impl` / `bf16_proj`; under `STRATA_GPU_METAL`, `Prefill::run` returns via `metal_engine::prompt_run` first (src/prefill/prefill.cpp:4814) |

Only `fused_gr_persist_init` (defect 3) is reached on the default path. A stub that stops the engine is the project's documented behavior for "never expected to be called here", so the 20 are correct as stubs; the one needs the no-op body.

## The fix (5 files, all under `metal/`)

Everything is inside the Metal-only project (`metal/` is a standalone CMake target, "never part of the CUDA / HIP / SYCL builds" — metal/CMakeLists.txt:1-2; the CUDA/HIP builds never include `metal/include`), so the PC builds are unaffected by construction. 58 insertions, 6 deletions:

- **`metal/include/metal_compat/cuda_runtime.h`** — add the six names, forwarding to `strata::gpu` like every other name in the file.
- **`metal/include/strata/gpu/runtime.hpp`** + **`metal/src/gpu/runtime.cpp`** — the six `strata::gpu` functions, each following the file's existing pattern:
  - `device_cc_major/minor`: `*v = 0; return fail(kErrorUnsupported)` — the same pattern as `device_sm_count` / `device_clock_khz` ("Apple reports no core count"). Effect: `device_is_sm75()` (src/core/verify.cpp:276) gets `ok = false` → `per_dev = 1` → `false`. Correct: an Apple GPU is not a Turing card, and `strata::sm75_default(0, 0)` is `false` anyway.
  - `cudaMemcpyPeer` → one-GPU: a plain device-to-device copy (`memcpy_sync`), `kErrorInvalidValue` for a non-zero device ordinal. Never actually run on a single-GPU Mac (`p2p_` stays false: `device_enable_peer_access` fails).
  - `cudaDeviceDisablePeerAccess` → `fail(kErrorUnsupported)`, the same as `device_enable_peer_access`; the caller clears the last error right after.
  - `stream_priority_range` → `(0, 0)`, `stream_create_with_priority` → `stream_create(s, f)`: one priority level, i.e. the default — matching the existing precedent in `mtp.cpp:438` ("HIP, Metal: the default priority"). Only reachable with `STRATA_PIPED_CONCURRENT=1` (opt-in), so default behavior is unchanged.
- **`metal/src/engine/stubs.cpp`** — regenerated with the project's own `metal/tools/gen_stubs.py` from the link log: the `bf16` stub updated to the current 10-parameter signature, the four `gr_*` stubs updated to their current signatures, the 16 remaining new entry points added (21 minus `fused_gr_persist_init`, which moves to the no-op below), plus the eight headers the regenerated set includes that the old file lacked.
- **`metal/src/engine/setup_host.cpp`** — `void fused_gr_persist_init() {}` next to the existing `fused_gr_check() {}` (defect 3); the corresponding line is *removed* from stubs.cpp to avoid a duplicate definition.

### Full patch

```diff
diff --git a/metal/include/metal_compat/cuda_runtime.h b/metal/include/metal_compat/cuda_runtime.h
index eb64243e..66c1937f 100644
--- a/metal/include/metal_compat/cuda_runtime.h
+++ b/metal/include/metal_compat/cuda_runtime.h
@@ -46,7 +46,9 @@ inline constexpr unsigned cudaEventDisableTiming = strata::gpu::kEventDisableTim
 inline constexpr unsigned cudaEventBlockingSync = strata::gpu::kEventBlockingSync;
 inline constexpr unsigned cudaStreamNonBlocking = strata::gpu::kStreamNonBlocking;
 enum cudaStreamCaptureMode { cudaStreamCaptureModeGlobal = 0, cudaStreamCaptureModeThreadLocal = 1, cudaStreamCaptureModeRelaxed = 2 };
-enum cudaDeviceAttr { cudaDevAttrClockRate = 13, cudaDevAttrMultiProcessorCount = 16, cudaDevAttrIntegrated = 18 };
+enum cudaDeviceAttr { cudaDevAttrClockRate = 13,
+                      cudaDevAttrComputeCapabilityMajor = 75, cudaDevAttrComputeCapabilityMinor = 81,
+                      cudaDevAttrMultiProcessorCount = 16, cudaDevAttrIntegrated = 18 };
 
 // ---- errors, devices
 inline cudaError_t cudaGetLastError() { return strata::gpu::get_last_error(); }
@@ -65,11 +67,14 @@ inline cudaError_t cudaDeviceGetAttribute(int* v, cudaDeviceAttr a, int d) {
         case cudaDevAttrIntegrated: return strata::gpu::device_integrated(v, d);
         case cudaDevAttrMultiProcessorCount: return strata::gpu::device_sm_count(v, d);
         case cudaDevAttrClockRate: return strata::gpu::device_clock_khz(v, d);
+        case cudaDevAttrComputeCapabilityMajor: return strata::gpu::device_cc_major(v, d);
+        case cudaDevAttrComputeCapabilityMinor: return strata::gpu::device_cc_minor(v, d);
     }
     return strata::gpu::compat_detail::bad();
 }
 inline cudaError_t cudaDeviceCanAccessPeer(int* v, int d, int p) { return strata::gpu::device_can_access_peer(v, d, p); }
 inline cudaError_t cudaDeviceEnablePeerAccess(int p, unsigned f) { return strata::gpu::device_enable_peer_access(p, f); }
+inline cudaError_t cudaDeviceDisablePeerAccess(int p, unsigned f = 0) { return strata::gpu::device_disable_peer_access(p, f); }
 
 // ---- memory and copies
 template <class T> inline cudaError_t cudaMalloc(T** p, size_t n) { return strata::gpu::malloc_device(p, n); }
@@ -93,10 +98,17 @@ inline cudaError_t cudaMemcpy2DAsync(void* d, size_t dp, const voi