Phase M4: matmul closeout -- sync tax measured + documented, offset-0 semantics verified, honest docs

Task 1 (per-call sync overhead): measured and decomposed, NOT optimized away --
torch exposes no safe handle to its MPS queue/stream, so the private-queue +
sync design stays, now with the cost precisely documented (MPS_STATUS.md 11.3).

- BNB_MPS_PROFILE=1 now decomposes each native matmul call into sched (commit
  -> GPU start), gpu (kernel), and done (GPU end -> wait return), on the
  mach_absolute_time timebase GPUStartTime uses. Steady-state fp16: the fixed
  tax is ~0.15ms/call (~0.02 encode + ~0.07 sched + ~0.07 done); the pre-call
  torch.mps.synchronize() itself is ~1us on an idle queue. That tax is ~52% of
  a 4096x4096 gemv call and ~10% of an M=512 gemm.
- Queue sharing rejected with evidence: torch 2.12.1 exposes no stream handle
  from Python; libtorch's at::mps::getCurrentMPSStream()/MPSStream internals
  are reachable only via dlsym'd C++ mangled symbols + header-derived ivar
  offsets (commandQueue()/queue() are inline, commit()/flush() private, encodes
  must run on torch's private _serialQueue) -- an unguardable ABI trap. Sharing
  only the queue would not even remove the pre-sync (command buffers execute in
  commit order; torch batches into an uncommitted buffer). Future directions
  (libtorch-linked extension, torch.mps.compile_shader) recorded with their
  costs in 11.3.
- New test_sync_discipline_interleave_stress pins the discipline: dependent
  torch writes interleaved with native matmuls on the same buffers, parity
  every iteration. Verified to have teeth: fails 30/30 with the pre-sync
  no-op'd, passes 100% with it.

Task 2 (offset-0 input copy): verified on this build, clone stays.
- A view's data_ptr() IS base + storage_offset*itemsize (raw arithmetic);
  objc-probing such an interior pointer SIGSEGVs uncatchably, so casting it
  would be silent corruption or a crash -- the clone in _ensure_native_buffer
  is load-bearing.
- Better than NEXT_MATMUL_PLAN 5 feared: untyped_storage().data_ptr() DOES
  recover the base MTLBuffer (passes the contract check), so offset binding is
  possible; not implemented because steady-state matmul inputs are offset-0
  (the clone ~never fires) and it would touch every ABI entry point for no
  measured benefit. Recipe recorded in 11.4; semantics pinned by
  test_view_data_ptr_is_base_plus_offset.

Task 3 (docs):
- MPS_STATUS.md: M2/M3 short notes consolidated into a proper 11 (fused gemv,
  native gemm, the M4 sync/offset findings); op matrix rows for gemv/gemm
  updated to native; header records the M4 harness result (328 passed + the
  known unrelated lion strict-xfail XPASS).
- docs/apple_silicon/README.md: gemv_4bit -> native (fused, 3.4-6.2x);
  gemm_4bit -> native for fp16/fp32 with the honest caveats (bf16 falls back --
  MPSMatrixMultiplication has no bf16 on macOS 26.4.1; large-M ~ parity with
  F.linear); Known limitations rewritten around the real remaining costs.
- README.md accelerator table, mps QLoRA 4-bit: 🐢 -> ✅ with a footnote.
  Justification: "slow implementation" is no longer accurate (inference gemv is
  fused native Metal at 3.4-6.2x; gemm is native fp16/fp32), but a bare ✅
  would overclaim, so the footnote carries the bf16-fallback and large-M-parity
  caveats inline in the README.
- "mis-cast" -> "miscast" (typos hook); _typos.toml learns the mach
  timebase field `numer` so the hook stops rewriting it to `number`.

Parity gate: BNB_MPS_REQUIRE_NATIVE=1 pytest tests/test_mps_parity.py ->
328 passed; the only failure is the pre-existing, unrelated lion strict-xfail
XPASS that fails identically on clean HEAD (separate optimizer track).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>

Author: eaglstun

README.md

@@ -182,12 +182,18 @@ bitsandbytes has the following minimum requirements for all platforms:
       <td>⬜ Metal <br><code>mps</code></td>
       <td>Apple M1+</td>
       <td>❌</td>
-      <td>🐢</td>
+      <td>✅ <sup>1</sup></td>
       <td>❌</td>
     </tr>
   </tbody>
 </table>
 
+<small><sup>1</sup> On <code>mps</code>, 4-bit matmul runs on native Metal kernels: a fused gemv for
+inference (M=1) and an <code>MPSMatrixMultiplication</code>-backed GEMM for fp16/fp32 batches. bf16
+batched matmul and other unsupported shapes use a slower dequantize+matmul fallback, and large-batch
+GEMM performs on par with dequantize+matmul. Details:
+<a href="./docs/apple_silicon/README.md">docs/apple_silicon</a>.</small>
+
 ## :book: Documentation
 
 - [Official Documentation](https://huggingface.co/docs/bitsandbytes/main)

_typos.toml

@@ -15,6 +15,7 @@ extend-ignore-re = [
 extend-ignore-identifiers-re = [
     ".*arange.*",
     ".*ARANGE.*",
+    "numer",  # mach_timebase_info_data_t.numer (csrc/mps_ops.mm)
 ]
 
 [type.py.extend-words]

csrc/mps_ops.mm

@@ -5,6 +5,7 @@
 #include <cstdint>
 #include <dlfcn.h>
 #include <libgen.h>
+#include <mach/mach_time.h>
 #include <string>
 
 // Native Metal dispatch layer for the bitsandbytes MPS backend.
@@ -133,6 +134,17 @@
     }
 }
 
+// Host time in seconds on the mach_absolute_time timebase -- the same clock
+// MTLCommandBuffer's GPUStartTime/GPUEndTime report, so the two are directly comparable.
+// Used only by the BNB_MPS_PROFILE probe.
+static double host_time_s() {
+    static mach_timebase_info_data_t tb = {0, 0};
+    if (tb.denom == 0) {
+        mach_timebase_info(&tb);
+    }
+    return (double)mach_absolute_time() * tb.numer / tb.denom / 1e9;
+}
+
 // One thread per block: cap the threadgroup and dispatch a non-uniform grid, then block on
 // completion. dispatchThreads is supported on all Apple Silicon GPUs.
 static void dispatch_per_block(id<MTLComputeCommandEncoder> enc, id<MTLComputePipelineState> pso, int64_t num_blocks) {
@@ -264,16 +276,26 @@ static void dispatch_per_block(id<MTLComputeCommandEncoder> enc, id<MTLComputePi
         // One SIMD-group per output element n.
         [enc dispatchThreadgroups:MTLSizeMake((NSUInteger)N, 1, 1) threadsPerThreadgroup:MTLSizeMake(32, 1, 1)];
         [enc endEncoding];
-        [cb commit];
-        [cb waitUntilCompleted];
 
-        // Kernel-only timing probe (BNB_MPS_PROFILE=1): wall-clock around this call includes
-        // the cross-queue sync + encode + blocking wait; GPUStartTime/GPUEndTime isolates the
-        // kernel itself for bandwidth accounting.
+        // Timing probe (BNB_MPS_PROFILE=1): decomposes the blocking call into
+        //   sched = commit -> GPU start (driver/queue scheduling latency)
+        //   gpu   = kernel execution (GPUStartTime..GPUEndTime)
+        //   done  = GPU end -> waitUntilCompleted return (completion delivery)
+        // Wall-clock around the whole ctypes call additionally includes encode (above) and
+        // the caller's torch.mps.synchronize().
         static const bool profile = getenv("BNB_MPS_PROFILE") != nullptr;
+        const double t_commit = profile ? host_time_s() : 0.0;
+        [cb commit];
+        [cb waitUntilCompleted];
         if (profile) {
+            const double t_done = host_time_s();
+            const double sched_ms = ([cb GPUStartTime] - t_commit) * 1000.0;
             const double gpu_ms = ([cb GPUEndTime] - [cb GPUStartTime]) * 1000.0;
-            NSLog(@"bnb_mps_gemv_4bit %@ N=%lld K=%lld gpu=%.3fms", name, (long long)N, (long long)K, gpu_ms);
+            const double done_ms = (t_done - [cb GPUEndTime]) * 1000.0;
+            NSLog(
+                @"bnb_mps_gemv_4bit %@ N=%lld K=%lld sched=%.3fms gpu=%.3fms done=%.3fms", name, (long long)N,
+                (long long)K, sched_ms, gpu_ms, done_ms
+            );
         }
     }
 }
@@ -427,15 +449,18 @@ static void dispatch_per_block(id<MTLComputeCommandEncoder> enc, id<MTLComputePi
             [enc endEncoding];
         }
 
+        static const bool profile = getenv("BNB_MPS_PROFILE") != nullptr;
+        const double t_commit = profile ? host_time_s() : 0.0;
         [cb commit];
         [cb waitUntilCompleted];
-
-        static const bool profile = getenv("BNB_MPS_PROFILE") != nullptr;
         if (profile) {
+            const double t_done = host_time_s();
+            const double sched_ms = ([cb GPUStartTime] - t_commit) * 1000.0;
             const double gpu_ms = ([cb GPUEndTime] - [cb GPUStartTime]) * 1000.0;
+            const double done_ms = (t_done - [cb GPUEndTime]) * 1000.0;
             NSLog(
-                @"bnb_mps_gemm_4bit dtype=%lld M=%lld N=%lld K=%lld bias=%d gpu=%.3fms", (long long)dtype_flag,
-                (long long)M, (long long)N, (long long)K, bias ? 1 : 0, gpu_ms
+                @"bnb_mps_gemm_4bit dtype=%lld M=%lld N=%lld K=%lld bias=%d sched=%.3fms gpu=%.3fms done=%.3fms",
+                (long long)dtype_flag, (long long)M, (long long)N, (long long)K, bias ? 1 : 0, sched_ms, gpu_ms, done_ms
             );
         }
     }