Skip to content

CHANGELOG.md

Latest commit

 

History

History
1339 lines (1208 loc) · 77.3 KB

CHANGELOG.md

File metadata and controls

1339 lines (1208 loc) · 77.3 KB
Raw

Changelog

Unreleased

Heterogeneous compute substrate validated end-to-end across two continents on four physically distinct machines, two GPU architectures (Apple Metal + NVIDIA CUDA tensor cores), two CPU ISAs (x86_64 + ARM), and a full DiLoCo training run with 528 bytes per outer-step over real TCP. The substrate has gone from "Apple Silicon GEMM kernels" to "vendor-neutral universal compute fabric for AI/HPC/graphics," with every architectural primitive in code and tested:

  • CPU-side attention/training/conv kernels (RMSnorm, LayerNorm, RoPE, SwiGLU, softmax, AdamW, FlashAttention-2 fwd+bwd, Conv2D fwd+bwd).
  • BLAS delegation (Accelerate / Intel MKL / OpenBLAS) — 2.10 TFLOPS fp32 at 4096³ on 88-core Xeon E5-2699 v4.
  • AVX2 / NEON / AMX in-tree GEMM micro-kernels (opt-in, self-contained).
  • AMX fp32 GEMM now overlaps X/Y tile loads with the previous FMA32 outer product in the inner loop while preserving the skip-Z first-FMA contract; direct AMX edge regressions remain bit-tight across transpose, alpha/beta, and K=0 cases.
  • Direct CUDA backend — tc_cuda_gemm validated against cuBLAS managed-memory dispatch on RTX 3090: 31.32 TFLOPS fp32, 32.28 TFLOPS fp16/fp32-accum, and 60.42 TFLOPS fp16-accum.
  • chipStar HIP backend scaffolding for Intel Level Zero, AMD OpenCL, ARM Mali — runs on every Khronos-standards GPU vendor.
  • HIP/chipStar fp32 GEMM dispatch now routes tc_gemm through TC_BACKEND_HIP / hipblas_sgemm_staged when the runtime initializes, with test_hip_gemm, scripts/ci_hip_smoke.sh, and JSON evidence validation for skipped vs passed HIP hosts.
  • GLOO TCP transport with full collective set: default brokered SUM/AVG/MIN/MAX, opt-in ring fp32 SUM/AVG for 3+ ranks, broadcast any-root, allgather, sparse_allreduce.
  • NAT-tolerant GLOO ring setup: direct neighbor links now use advertised per-rank addresses, bounded connect timeouts, and coordinated broker fallback when any ring edge is unreachable.
  • GLOO ring observability: TC_GLOO_TRACE=1 logs direct-ring enablement, coordinated fallback, and fp32 SUM route selection so WAN runs prove whether they used the ring or broker.
  • Distributed runtime evidence now claims GLOO teardown explicitly: run_distributed_runtime_evidence.py covers close_gloo_state when the forked ring smoke passes, so setup and cleanup are both visible to ICC.
  • Metal ops evidence now includes the batched-GEMM host validation path: run_metal_ops_runtime_evidence.py runs test_gemm_f16 and covers lib/ops/gemm.mm:batched_matrix_bytes.
  • Metal ops evidence now also runs test_fused_norm_gemv, covering the tg_sum32 shader reduction helper, and GEMM tracing reports selected Metal kernel names so async-copy wrapper coverage is claimed only when an async-copy kernel is actually selected at runtime.
  • TensorOps selection now includes the planned narrow Metal 4 attention kernel/shape envelope (D=64/128, tiled sequence lengths, no LSE/window/ ALiBi variants), with unit coverage before the M5 runtime path is promoted.
  • CPU ops evidence now claims the Conv2D shape guard and im2col staging helper when test_conv2d passes, covering conv_dims_valid and im2col_fp16.
  • Python packaging evidence now covers the full setup.py native-artifact helper chain, including artifact discovery, platform selection, metallib requirement checks, macOS dylib/tag validation, and _run_tool.
  • AMX/bench evidence now covers benchmark parser/timing helpers in bench_gemm.c and can cover the optional FlashAttention helpers in bench_attention.c when the host exposes Metal to evidence subprocesses.
  • AMX/bench evidence now claims the AMX worker-pool one-time initializer when the two-worker AMX GEMM path dispatches successfully.
  • AMX/bench evidence now keeps FlashAttention benchmark coverage optional when nested Python evidence subprocesses cannot acquire a Metal device, preventing AMX worker evidence from being blocked by host Metal visibility.
  • GeoRefine Qwen rank-probe launches now pass --worker-lease-mode mirror through the qLLM lease wrapper so scheduler authority and worker lease handoff stay aligned.
  • CUDA smoke evidence now names the shared block_reduce_sum_f32 helper, and configs/tensorcore_cuda_smoke_job.template.json provides a scheduler-owned remote CUDA kernel proof path for the remaining NVIDIA-only evidence.
  • The CUDA smoke submission template now takes admission helper-process allowances from inventory reconciliation settings instead of embedding them in the checked-in template.
  • test_dist_remote now accepts --elements and --iters, making short WAN ring proofs possible without changing the default 4 MB throughput probe.
  • Fused quantized RMSNorm GEMV is now closed across the public release surface: shared-library exports, Python ctypes binding, and release-smoke source-manifest coverage.
  • CUDA/portable training closure now registers the end-to-end training convergence test and native training_step example in non-Metal CTest, so CUDA hosts prove more than isolated training kernels.
  • CUDA smoke now emits machine-readable runtime evidence and proves the Python training dispatch path, including RMSNorm/SwiGLU/softmax backward plus fp32/fp16-gradient AdamW, not just cuBLAS GEMM.
  • CUDA training dispatch now covers LayerNorm backward and RoPE forward/backward on managed buffers, closing two more transformer-training fallbacks in the RTX 3090 path and adding them to CUDA smoke evidence.
  • Windows x86 portable CPU bring-up now has MSVC-safe CMake flags, a generated DLL export definition from the public ABI list, Windows-safe GGUF mapping/autotune helpers, and scripts/ci_windows_cpu.ps1 for Jack's Tailscale machine and hosted Windows CI.
  • Portable CPU GEMM now honors the BLAS beta == 0 contract without reading stale C contents, fixing the Windows conv2d forward NaN regression and adding poisoned-output coverage for f32/f16/bf16/i8.
  • Jack's Windows host now has a reproducible bootstrap path: scripts/bootstrap_windows_cpu.ps1 detects Visual Studio Build Tools, CMake/CTest, and Python, then runs the full Windows portable CPU smoke. MinGW-style and MSVC-style DLL names are both accepted by the loader and packaging scripts.
  • Windows host smoke is now operational from the Unix side too: scripts/run_windows_host_smoke.sh SSHes over Tailscale, clones or fast-forwards the remote checkout, and runs the same bootstrap gate on Jack's machine or any future Windows node. The target SSH coordinate now comes from TC_WINDOWS_SSH or a private local config file instead of a source-controlled tailnet address.
  • Windows host smoke can now emit machine-readable evidence, with scripts/check_windows_host_smoke_evidence.py and operational-bundle policy flags for requiring clean-head Windows proof.
  • Windows GLOO TCP is now backed by Winsock and covered by a local split-rank CTest launcher, so Jack's x86 Windows host proves TC_DIST_GLOO allreduce over loopback instead of only portable CPU single-process kernels.
  • Eshkol bridge diagnostics now expose C-string device/backend/version/status pointers through real tc_eshkol_* C shims, exported with the existing bridge ABI and exercised by the Scheme bridge smoke.
  • Eshkol bridge smoke now defaults to build-portable-cpu-current when that build exists, giving backend-independent --require-pass evidence for the Scheme/C bridge while preserving --build-dir build for Metal-path probes.
  • Metal kernel builds now track local Metal headers, emit collision-resistant AIR filenames, and fail early with explicit xcrun metal/metallib discovery diagnostics.
  • Mesh resource scheduling now supports completion gates: GeoRefine Qwen jobs can declare an artifact checker, completed jobs are not relaunched, completed stale leases are released, and unknown completion checks block relaunch for that scheduler pass.
  • Mesh resource scheduling now supports admission gates: jobs can declare an admission_cmd over CUDA/HIP/Windows/live-mesh evidence, and failed or timed-out admission blocks new launches without disturbing already-live holders.
  • Mesh resource scheduling now has a cuda_exclusive resource class with required admission, post-start, and worker-identity gates, plus scripts/check_cuda_resource_admission.py and scripts/mesh_worker_identity.py for worker host process, cgroup, and CUDA metadata.
  • Mesh resource scheduling now has a checked-in fleet inventory (configs/mesh_resources.json) plus validation and arbiter-capacity wrapper tools, so blocked hosts and reserved slots such as Enki's Tsotchke-chan M4 lane are rejected before any lease is claimed.
  • Mesh inventory backend: "cuda" rows now force cuda_exclusive scheduler semantics, preventing jobs from downgrading CUDA accelerators to generic resources and bypassing admission/post-start/worker-identity gates.
  • Mesh resource scheduling now supports distributed pool placement: resources and resource_pool jobs expand across inventory resources, command templates receive {resource}/{node} context, and tenant-aware fair-share selection prevents one user from monopolizing every idle machine in a single scheduler pass.
  • Mesh resource scheduling now has a tensorcore.job.v1 submission boundary with submit/status/cancel/drain/undrain/audit commands plus a scheduler-owned GeoRefine Qwen CR070 rank-probe template. The template pins run identity through liveness, worker-identity, and completion gates, and cancel accepts the expanded placement ids reported by status.
  • Mesh scheduler queue mutations now require an explicit append-only --event-log-jsonl, CUDA inventory rows declare per-resource GPU reconciliation policy, and scripts/mesh_worker_gpu_reconcile_sweep.py sweeps worker GPU occupancy without hardcoded arbiter or resource aliases. scripts/mesh_gpu_reconciliation_audit.py wraps the sweep plus scheduler audit for one-step control-plane checks, and the scheduler can require that fresh audit artifact before placing CUDA jobs. Scheduler loop/status/audit commands can also require queue event-log integrity so out-of-band queue edits fail closed.
  • The GeoRefine Qwen rank-probe launcher now takes host, path, model, worker resource, and quality gate values from scheduler config or environment instead of embedding cosbox/private-path defaults in code.
  • scripts/check_mesh_resource_jobs.py now rejects rendered GeoRefine submission templates that pin concrete hosts, worker aliases, private paths, or queue state instead of using selector/placeholders plus TC_GEOREFINE_*.
  • Mesh scheduler heartbeats for live CUDA-exclusive jobs now refresh worker identity metadata on the existing lease instead of leaving identity pending.
  • Mesh scheduler jobs now have a checked-in source of truth (configs/mesh_resource_jobs.json) plus scripts/check_mesh_resource_jobs.py, and live jobs can safely adopt same-tenant manual leases when explicitly declared metadata fields match.
  • scripts/mesh_system_audit.py now audits the whole mesh control plane: inventory, expanded scheduler jobs, scheduler freshness, arbiter resources, reserved/blocked policy, scheduler lease identity metadata, and optional live CUDA process ownership.
  • Jack's Windows host is now proven over Tailscale/SSH for portable CPU bootstrap, CTest, install smoke, and Python smoke on the current head.
  • Windows CUDA readiness now has a non-destructive probe and validator: scripts/run_windows_cuda_probe.sh records nvidia-smi device facts, compute-app admission state, CUDA Toolkit / nvcc discovery, and clean git-head provenance for Jack's RTX lane.
  • Jack's RTX 3060 now has current driver, admission, CUDA Toolkit 12.6 redistributable, Windows CUDA configure/build, full CTest, and CUDA GEMM runtime evidence.
  • Windows CUDA probing now distinguishes opaque WDDM desktop rows from real CUDA process-table entries and can discover Jack's user-local CUDA 12.6 redistributable toolkit under src/cuda-redist-12.6; optional build-smoke evidence now records CUDA configure/build/CTest and test_cuda_gemm.
  • Jack's RTX 3060 lane is now active in the checked-in scheduler inventory as a paused CUDA-exclusive lane with WDDM-aware admission and Windows worker-identity helpers. It cannot launch work until a submitted job provides real start and post-start probes for the specific Windows CUDA process.
  • Mesh deployment is now git-checkout based: scripts/mesh_deploy_git_checkout.py clones or fast-forwards repos on SSH nodes, checked-in scheduler configs reject private wrapper paths, Jack's scheduled CUDA smoke uses repo-local helpers but stays paused until Jack has a persistent Windows launch path, and GeoRefine CR025 and old-donkey qLLM precompute now have repo-owned remote starters with non-launching preflight modes. GitHub SSH access is configured on the target hosts; the rows stay paused for operator-controlled adoption or launch. Default preflight sweeps now include GeoRefine, old-donkey, and Jack so paused launchable rows stay visible, while any future opt-out rows remain explicit via skipped_default_job_ids and --mesh-preflight-skipped-default-job. Windows SSH helpers now upload large probe scripts with scp instead of stdin-fed PowerShell to avoid Jack's OpenSSH stdin-close behavior. The stale qLLM phase-1 systemd row is adoption-only until its launcher is installed from a qLLM git checkout.
  • Release-smoke evidence now records clean git-head provenance, and the operational bundle checker can require release, SDK26, PyTorch, and live-mesh evidence to match the current committed head.
  • Metallib build-rule evidence now probes a generated CMake project that calls tc_compile_metallib, hashes the generated .metallib, emits ICC-readable coverage for the CMake helper, and reports explicit blocked statuses on hosts without Apple Metal tools. The CMake helper now also emits an explicit runtime backend probe line for xcrun metal/metallib discovery, so ICC can distinguish observable Metal-tool gating from an unresolved external backend.
  • Python packaging evidence now directly proves the native artifact copy path and macOS validation-tool path in setup.py, including copied artifact hashes, wheel hash, explicit platform tag, and ICC-readable coverage for _run_tool, build_py_with_native_artifacts.run, and bdist_wheel_with_native_artifacts.run. _run_tool now has opt-in structured probe logging, and the evidence runner records that log without polluting parsed lipo architecture results.
  • Distributed runtime evidence now wraps the forked GLOO ring, dense DiLoCo, and sparse DiLoCo smokes, recording explicit loopback-blocked states when local TCP sockets are unavailable and ICC-readable coverage for the GLOO ring connection helpers plus DiLoCo outer-step and TOPK sparse-delta core.
  • AMX and GEMM benchmark evidence now wraps the AMX metadata probe, direct AMX GEMM regression, and a tiny bench_gemm run, emitting ICC-readable coverage for AMX worker/core entry points plus bench_one while recording TensorOps M5 layout helpers as optional skipped_no_metal4_sdk/skipped_no_m5 blockers when the host cannot genuinely execute them.
  • Portable CPU ops evidence now wraps test_portable_cpu and test_conv2d, emitting ICC-readable coverage for gemm_compute, CBLAS f16/bf16 helpers when compiled, fp16 conversion, and the Conv2D validation/backward helpers.
  • Metal ops evidence now wraps attention correctness and Metal Conv2D smokes with TC_TRACE=1, emitting coverage for attention encode_forward and Conv2D conv_bytes while leaving shader-internal async-copy coverage explicitly blocked until selected-kernel or shader-line instrumentation exists.
  • Quantized/GGUF runtime evidence now wraps test_quantized and test_gguf, emitting ICC-readable coverage for the Metal gemv_quant_encode helper and the GGUF quantized-matrix descriptor helper while reporting explicit metal_device_unavailable blocked states when the host sandbox hides the Metal device.
  • Eshkol bridge evidence now normalizes expected no-GPU command tails in skipped_no_gpu artifacts, preserving hashes while preventing generic ICC failure scanning from treating a deliberate hardware skip as a runtime failure.
  • Hardware evidence dispatches now include a GitHub-hosted runner preflight artifact and CI-tested preflight script so missing [self-hosted, macOS, ARM64] runners are visible before the hardware job queues. A live dispatch validated that the preflight artifact uploads even when the self-hosted Apple GPU job remains queued. M5 TensorOps runners now also have a local host preflight that records macOS/arm64, Xcode/SDK26, display device, and optional test_tensorops_runtime readiness before the full hardware smoke runs. M5 preflight artifacts now include git provenance, structured environment/source diagnostic classes, a dedicated checker, and operational-evidence wiring; a fetch helper can dispatch or download the required self-hosted M5 workflow artifact and validate it against a clean expected head. The GitHub-hosted hardware runner preflight artifact now also records workflow/run/head provenance and has a checker for validating runner visibility diagnostics after download. Its diagnostics now include the explicit next action for token, absent-runner, offline-runner, and online-runner states. The M5 evidence fetch helper can now fetch that preflight artifact for queued runs and cancel the workflow automatically when no online matching runner is visible. The required M5 TensorOps workflow path now routes only to explicitly labelled [self-hosted, macOS, ARM64, m5, sdk26, metal4-tensorops] runners, the preflight checker enforces those labels for Metal 4 TensorOps dispatches, and the fetch helper no longer auto-cancels API/token failures as if they were known offline-runner states.
  • HIP smoke evidence now proves both fp32 hip_gemm_sgemm and fp16 hip_gemm_hgemm, records per-kernel gemm_kernels, and emits ICC-readable source coverage for remote chipStar/HIP runs.
  • HIP diagnostic evidence now records readiness.diagnostic_class, separating clean no-HIP/ROCm absence from an indicated but blocked HIP setup, and both HIP evidence checkers can require that class with clean-head provenance.
  • M5 TensorOps preflight candidate evidence now requires proven host, Xcode, SDK 26, and display-GPU checks; unverifiable display probes remain blocked instead of becoming candidates.
  • CUDA smoke evidence now records per-kernel fp32/fp16/bf16/int8 GEMM status, AdamW dispatch coverage, and ICC-readable source coverage for remote NVIDIA runs that support those paths.
  • Eshkol bridge runtime evidence now has a repo-owned smoke/checker pair: scripts/run_eshkol_tensorcore_bridge_smoke.py records the real local eshkol-run compile/runtime state for hello_tensorcore.esk and the new broader tensorcore_bridge_smoke.esk. The Scheme bindings now resolve __tc-* through repo-owned tc_eshkol_* C shims, so the bridge smoke can pass end-to-end against the portable CPU backend while still failing cleanly on hosts where the selected native backend cannot initialize.
  • CUDA smoke evidence now participates in the same clean-head gate, including archive-based remote builds prepared by the live mesh runner.
  • HIP smoke evidence now uses the same source-head override/marker mechanism and operational clean-head gate as CUDA.
  • HIP evidence validation now has a middle policy gate: --require-hip-build requires that chipStar/HIP runtime targets compiled, while still accepting explicit runtime-unavailable evidence on hosts such as NVIDIA/OpenCL where SPIR-V device initialization is unavailable.
  • HIP/OpenCL/SPIR-V readiness now has a standalone probe: scripts/probe_hip_toolchain.py records hipcc, LLVM/SPIR-V translator, OpenCL/Level Zero runtime, GPU/SPIR-V device capability, CMake package, hipBLAS, and path-hint evidence; scripts/check_hip_toolchain_evidence.py and the operational bundle can require build-toolchain, SPIR-V-capable GPU runtime, or ready-for-hipBLAS policy before a host is admitted as a chipStar accelerator. Archive-based remote runs now honor .tensorcore_source_head / .tensorcore_source_dirty provenance markers, and the probe recognizes versioned LLVM/SPIR-V tools such as llvm-spirv-19.
  • DiLoCo runtime with NONE/FP16/TOPK_1PCT/TOPK_01PCT compression, SGD/Nesterov/Adam outer optimizers, async overlap, sparse-on-the-wire cross-continent path.
  • Memory-tier API (L0-L4) and activation checkpointing API.
  • Activation checkpointing now serializes same-id discard/realize/unregister operations with per-checkpoint locks and atomic counters, with a dedicated concurrency CTest covering one-callback realize fan-out and one-accounted discard fan-out.
  • Cross-process test infrastructure (test_diloco_gloo_fork, test_diloco_sparse_fork, test_dist_remote) for forked, multi-rank, and cross-machine validation.
  • examples/mesh_training_demo.c: runnable split-rank training loop with RMSNorm, GEMM, softmax+CE, AdamW, and DiLoCo outer synchronization; registered as a single-rank CTest smoke and parameterized for GLOO multi-host rendezvous.
  • scripts/run_live_mesh_training_demo.sh: one-command live mesh training runner for Atlas + Enki + old-donkey + cosbox. The runner can prepare remote binaries from the current committed checkout, build cosbox with TC_ENABLE_CUDA=ON, launch the four ranks over Tailscale, and verify DiLoCo outer sync plus checkpoint counters on every rank. It can also write machine-readable live-run evidence, validated by scripts/check_live_mesh_training_evidence.py.
  • Live mesh training evidence now records requested vs observed backends per rank and fails by default when rank 3 is configured for CUDA but does not report backend=cuda; intentional fallback runs must set TC_MESH_ALLOW_CUDA_FALLBACK=1 and are visible in the JSON summary.
  • Add localhost mesh-training evidence mode: set TC_MESH_LOCAL_ONLY=1 to run all mesh_training_demo ranks on one host with the same direct ring, DiLoCo, checkpoint, and evidence parser path. The evidence checker now has --require-local-only for regression gates when the physical mesh is intermittent.
  • Add rank-indexed GLOO ring advertise hosts: TC_GLOO_ADVERTISE_HOSTS accepts a comma-separated rank list for direct-ring neighbor dialing on Tailscale/NAT overlays, while preserving coordinated broker fallback when any ring edge is unreachable.
  • Propagate GLOO advertised hosts in the live mesh launcher: scripts/run_live_mesh_training_demo.sh now passes TC_GLOO_ADVERTISE_HOSTS through to all local and remote ranks.
  • Add per-rank remote PATH overrides for live mesh training: TC_MESH_RANK{1,2,3}_PATH lets operators prepend host-specific toolchain directories during remote prepare and rank launch.
  • Allow rank 1 Linux prepare in the live mesh launcher: TC_MESH_RANK1_PREPARE=linux builds Enki/rank 1 from the archived checkout like the Linux ranks instead of copying the local Apple binary.
  • Record live mesh launch topology in evidence: live training evidence now includes per-rank launch/prepare metadata, and the checker can require rank 1 source preparation with --require-rank1-source-prepare.
  • Move live mesh coordinates to private config: run_live_mesh_smoke.sh and run_live_mesh_training_demo.sh now read TC_MESH_CONFIG or explicit TC_MESH_RANK* environment variables instead of carrying source defaults for private hostnames or tailnet addresses.
  • Add operational evidence bundle validation: scripts/check_operational_evidence.py validates release, SDK26, CUDA, HIP, PyTorch, mesh preflight, and live-mesh artifacts together, with clean-current-head enforcement for physical mesh deployment evidence.
  • Mesh training activation checkpoint mode: mesh_training_demo --checkpoint now discards X_norm after the forward projection, realizes it through the RMSNorm recompute callback before the dW backward GEMM, and has a dedicated CTest smoke.
  • Experimental PyTorch bridge with zero-copy fp32 CPU matmul, an opt-in torch.matmul dispatcher hook, host-memory PrivateUse1 allocation, explicit to_tensorcore() / to_cpu() round-trips, and structured backend-state reporting. The bridge now also exposes matmul_eligibility() so training code can see the exact tensorcore-dispatch versus ATen-fallback reason.

End-to-end validation matrix

Run Hardware Result
test_gloo_ring_fork (forked) Atlas M2 Ultra ✓ — opt-in 4-rank TCP ring SUM
test_diloco_gloo_fork (forked) Atlas M2 Ultra ✓ — 2 ranks converge
test_diloco_sparse_fork (forked) Atlas M2 Ultra ✓ — 16× bandwidth reduction
Atlas ↔ Enki (Tailscale) M2 Ultra + M4 ✓ — cross-arch Apple
Atlas ↔ old-donkey (Tailscale, cross-continent) Mac + Linux Xeon ✓ — 1.4 s for 3 outer steps
Atlas ↔ cosbox (Tailscale, CUDA-built) Mac + RTX 3090 host ✓ — TC_ENABLE_CUDA validated
4-rank cross-continent: Atlas + Enki + old-donkey + cosbox Mac×2 + Linux×2, two continents ✓ — first 4-way mesh
mesh_training_demo live 4-rank Atlas + Enki + old-donkey + cosbox ✓ — direct-ring DiLoCo outer sync + activation checkpointing; 5 outer-step soak completed; cosbox rank 3 used CUDA
tc_cuda_init device introspection RTX 3090 Ampere sm_8.6 ✓ — fp16+bf16+int8_tc+tf32
tc_gemm via cuBLAS tensor cores RTX 3090 4096³ fp16 ✓ — 32.28 TFLOPS fp32-accum / 60.42 TFLOPS fp16-accum
tc_gemm via cuBLAS sgemm RTX 3090 4096³ fp32 ✓ — 31.32 TFLOPS
CUDA training CTest + Python smoke cosbox RTX 3090 at 6382b98 ✓ — 18/18 CTest; RMSNorm/SwiGLU/softmax backward + AdamW CUDA dispatch

Observability

  • Widen backend diagnostics beyond GEMM: tc_last_backend() now records served training, Conv2D, quantized, attention, GEMM, TensorOps, and portable CPU dispatches. Generic Metal kernels report TC_BACKEND_METAL_COMPUTE = 8 / "metal_compute".
  • Add TC_TRACE=1 dispatch logs: served compute dispatches can emit op/status/backend lines to stderr, with portable CPU smoke coverage checking GEMM and softmax traces.

CPU compute stack expanded

  • Add CPU FlashAttention forward + backward: lib/ops/attention_cpu.cpp. Memory-efficient online-softmax algorithm (Br=Bc=32 tiles), GQA + causal + sliding window + ALiBi all supported, OpenMP per-(B, H) parallelism. fp16 IO, fp32 accumulator.
  • Add CPU training kernels: lib/ops/training_cpu.cpp implements tc_rmsnorm_forward/backward, tc_layernorm_forward/backward, tc_rope_forward/backward, tc_swiglu_forward/backward, tc_softmax_forward/backward, tc_adamw_step (fp16 and fp32 grad paths), tc_fused_rmsnorm_gemv, and tc_fused_layernorm_gemv. All OpenMP-parallel.
  • Add fused LayerNorm+GEMV public primitive: tc_fused_layernorm_gemv mirrors the RMSNorm decode projection fast path for LayerNorm-based models, with Metal, portable CPU, Python, export-surface, and separate tc_layernorm_forward + tc_gemm correctness coverage.
  • Add CPU Conv2D forward + backward: lib/ops/conv2d_cpu.cpp via im2col/col2im + GEMM, inheriting the BLAS-delegate fast path and matching the Metal path's shape/buffer validation.
  • Wire BLAS-delegate (Accelerate / MKL / OpenBLAS) into CPU GEMM: lib/ops/gemm_cpu.cpp detects CBLAS at CMake time and delegates fp32 / fp16-through-fp32 GEMM. Measured on old-donkey (Xeon E5-2699 v4 x 2 sockets):
    • reference triple-loop: 0.66 GFLOPS at 1024³
    • OpenBLAS 44-thread: 1.34 TFLOPS at 4096³
    • Intel MKL 44-thread: 2.10 TFLOPS at 4096³
  • Harden the PyTorch bridge smoke: scripts/ci_pytorch_smoke.sh now force-builds bindings/pytorch, validates fp32, bf16, non-contiguous tensors, K == 0 / empty-result matmuls, error paths, and the opt-in torch.matmul dispatcher/autograd fallback. The bridge now handles degenerate PyTorch matmuls at the binding boundary instead of passing zero-byte buffers into the C ABI. The smoke also pre-initializes the ctypes C ABI before importing the extension, covering the TC_ERR_ALREADY_INITIALIZED path that appears in mixed Python integrations.
  • Emit PyTorch bridge smoke evidence: scripts/ci_pytorch_smoke.sh writes optional JSON via TENSORCORE_PYTORCH_SMOKE_EVIDENCE_PATH, and scripts/check_pytorch_smoke_evidence.py validates pass/skip state, backend registration, matmul dispatch coverage, and direct device="tensorcore" allocation status for node-health automation.
  • Harden distributed collective validation: Apple and portable GLOO collectives now check byte-count and allgather total-size overflow before buffer validation, and the hidden GLOO transport checks its own send/recv byte counts. The portable CPU suite now regresses allreduce, broadcast, and allgather overflow rejection.
  • Add hand-tuned AVX2 fp32 GEMM micro-kernel: lib/ops/gemm_cpu_avx2.cpp. 6×16 BLIS-style inner kernel, 12 ymm accumulators, FMA inner loop, shared A/B panel packing, and OpenMP fanout over independent tile work. Opt-in via TC_USE_AVX2_GEMM=1; TC_AVX2_THREADS=1 forces serial A/B runs. Self-contained and hidden from the public export surface.
  • Prove AVX2 GEMM 1024+ scaling: bench_gemm_shared now links the OpenMP-enabled shared runtime, while bench_gemm keeps measuring the dependency-light static SDK path. On old-donkey, TC_USE_AVX2_GEMM=1 TC_AVX2_THREADS=64 build/bench/bench_gemm_shared reaches 0.74 TFLOPS at 4096³ versus 1.53 TFLOPS for the MKL default.
  • Add hand-tuned NEON fp32 GEMM micro-kernel: lib/ops/gemm_cpu_neon.cpp. 8×8 aarch64 SIMD kernel for Apple/ARM CPU builds, opt-in via TC_USE_NEON_GEMM=1, with CBLAS remaining the default until broader throughput data is collected.
  • Normalize K==0 portable GEMM: the CPU path now treats the degenerate matrix product as C := beta*C for f32, f16, bf16, and i32 outputs, preserves padded ldc regions, and covers that behavior in test_portable_cpu.
  • Add opt-in Apple AMX fp32 GEMM prototype: lib/ops/gemm_cpu_amx.cpp. Apple-Silicon-only 16×16 fp32 tile path, gated by TC_USE_AMX_GEMM=1 and falling through to NEON/CBLAS when disabled. The committed path remains opt-in, handles transpose A/B, M/N edge tiles, K==0, and alpha/beta through a pad-and-trim wrapper, and uses persistent pthread worker-local packs for M>=256 only when tc_amx_cluster_count() reports two AMX-capable P-clusters. TC_AMX_THREADS=1 preserves the single-worker path for A/B checks. The worker pool is guarded for concurrent callers and reports worker allocation failures back to the GEMM fallback path. The FMA loop now uses AMX FMA32 skip-Z instead of zero-loading each tile, tc_amx_isa_version() records AMX1/2/3 from hw.cpufamily, and f16/bf16 entry points remain gated until FMA16 IO-mode hardware validation. Direct test_amx_gemm / test_amx_edge regressions compile in portable builds but only execute when TC_RUN_AMX_GEMM_TEST=1 is set, because hosted macOS runners can trap reverse-engineered AMX instructions. test_amx_probe runs unconditionally for the safe metadata/stub contract. CBLAS remains the default path.

Heterogeneous-mesh substrate

  • Add DiLoCo public ABI and local runtime: include/tensorcore/diloco.h, lib/distributed/diloco.cpp, and docs/diloco.md. The single-rank path is implemented and covered in portable CPU tests; dense multi-rank outer steps over TC_DIST_GLOO are covered by a forked localhost smoke. Sparse TOPK outer steps now use GLOO sparse packed all-reduce and have a separate forked localhost smoke; dropout-tolerant WAN recovery remains staged.
  • Add HIP/chipStar backend scaffolding: include/tensorcore/hip.h, lib/hip/device.cpp, lib/hip/gemm.cpp, and lib/hip/README.md. The public API exports deterministic unsupported behavior when HIP is not compiled in, while the split TUs stage chipStar device discovery and hipBLAS dispatch for Intel Level Zero plus NVIDIA/AMD/ARM OpenCL.
  • Add direct CUDA backend scaffolding: include/tensorcore/cuda.h, lib/cuda/device.cpp, and lib/cuda/gemm.cpp expose NVIDIA-native device diagnostics, hidden managed-buffer hooks, and a hidden cuBLAS dispatch hook. TC_BACKEND_CUDA now identifies successful default cuBLAS GEMM dispatches after CUDA initialization; runtime allocations use CUDA managed memory in that mode, while wrapped host pointers use the staged fallback. TC_DISABLE_CUDA_GEMM=1, TC_CUDA_GEMM=0, or TC_USE_CUDA_GEMM=0 force CPU fallback for A/B testing. scripts/ci_cuda_smoke.sh and test_cuda_gemm cover fp32/fp16 GEMM on CUDA hosts, including a 4096^3 managed-memory perf gate on high-end Ampere+ devices. CUDA builds also route bf16/fp32-accum and int8/i32-accum GEMM through cuBLAS when the device reports support. CUDA builds also compile managed-memory training kernels for RMSNorm forward/backward, LayerNorm forward/backward, RoPE forward/backward, SwiGLU forward/backward, softmax forward/backward, and fp32/fp16-gradient AdamW with CPU fallback for host-only buffers. test_training_kernels now requires CUDA dispatch for these managed-memory paths when a CUDA device is visible and CUDA has not been explicitly disabled. Default builds return deterministic unsupported statuses until TC_ENABLE_CUDA is wired to a CUDA toolchain.
  • Extend async Metal GEMM to bf16: kernels/metal/gemm_async.metal now has a templated fp16/bf16 async-copy body and exports tc_gemm_bf16_f32_async; dispatch probes it on Apple9+ when the async simdgroup path is forced over the MPS fallback.
  • Add opt-in CUDA/HIP CMake detection: TC_ENABLE_CUDA=ON now enables the direct CUDA scaffolding only when CMake finds CUDA::cudart and CUDA::cublas; TC_ENABLE_HIP=ON requires HIP runtime plus hipBLAS imported targets. Installed CMake packages rediscover those dependencies before loading tensorcore targets.
  • Wire portable CPU GLOO TCP collectives: lib/distributed/gloo_tcp.cpp now backs public multi-rank TC_DIST_GLOO on the portable CPU build for fp32 SUM/AVG/MIN/MAX all-reduce, fp16 SUM/AVG all-reduce, byte-level broadcast, allgather, and barrier. The localhost fork smoke validates the public path end-to-end.
  • Wire Apple GLOO TCP collectives: default Metal-enabled builds now route TC_DIST_GLOO through the same TCP transport as the portable CPU backend. The default CTest suite registers the GLOO, DiLoCo-over-GLOO, and sparse TOPK-over-GLOO fork smokes so Darwin distributed behavior is exercised before release.
  • Add GLOO TCP ring all-reduce: TC_DIST_GLOO can build direct ring neighbor sockets for world_size >= 3 and route fp32 SUM through a reduce-scatter/all-gather ring when TC_GLOO_RING=1 is set. The broker path remains the default for NAT-hostile multi-host deployments, while test_gloo_ring_fork covers the ring path in default and portable CTest.
  • Add opt-in PyTorch matmul dispatcher hook: bindings/pytorch can now route eligible fp32/bf16 CPU torch.matmul calls through tensorcore when tensorcore_torch.set_default_matmul(True) is enabled. The bridge uses zero-copy buffer wrappers when alignment permits and staged buffers otherwise, while unsupported shapes and dtypes fall back to ATen.
  • Add memory-tier public ABI: include/tensorcore/memory_tier.h and lib/core/memory_tier_stub.cpp expose buffer tier hints, async promote/demote entry points, and usage accounting. The shipped baseline is intentionally L0-only until L1-L4 hosting lands.
  • Finish zero-copy host-buffer wrapping: tc_buffer_from_ptr is now in the exported ABI, bound in Python, documented, and covered by portable C/Python smokes. Metal builds now expose the same map/free wrapper contract for page-aligned no-copy MTLBuffer views.
  • Implement portable CPU activation checkpointing: tc_checkpoint_discard now frees owned CPU buffer storage while preserving the handle, tc_checkpoint_realize reallocates and invokes the registered recompute callback, and test_checkpoint validates the lifecycle.
  • Add Metal activation checkpoint storage detach: Metal buffers now release and recreate the underlying MTLBuffer while preserving the public tc_buffer handle, so test_checkpoint runs in the default Apple suite instead of skipping.
  • Document cross-continent training topology: docs/diloco.md explains the DiLoCo algorithm, compression choices, and the two-site bandwidth budgeting model.
  • Expose HIP and DiLoCo through Python: the binding now has public hip_* wrappers, diloco_* wrappers, DiLoCoContext, enum constants, ABI-layout checks for the new structs, and portable-CPU smoke coverage for HIP inactive diagnostics plus single-rank DiLoCo outer steps.
  • Expose memory-tier controls through Python: buffer_set_tier_hint, buffer_get_tier, buffer_promote_async, buffer_demote_async, buffer_tier_sync, and memory_tier_usage are bound and covered by the portable-CPU smoke.
  • Expose activation-checkpointing controls through Python: the binding now covers checkpoint_* lifecycle and counter functions, with portable C and Python smoke coverage.

Original v0.1.22 changes (preserved below)

Portable CPU backend for non-Apple workers, documentation overhaul, executable examples for the inference + training assembly, hardened shell scripts, expanded CI breadth, and a packaged native SDK artifact on top of v0.1.22:

Portable CPU backend (TC_ENABLE_METAL=OFF)

  • Add portable CPU backend — lib/core/device_cpu.cpp, lib/ops/gemm_cpu.cpp, lib/ops/quantized_cpu.cpp, lib/ops/unsupported_cpu.cpp, lib/core/cpu_float.h. Pure C/C++17; builds on Linux, Intel Mac, or anywhere with a C++17 toolchain. New TC_BACKEND_PORTABLE_CPU = 7 enum value; tc_backend_name renders it as "portable_cpu".
  • Tighten portable CPU collectives — tc_dist_* works in the TC_DIST_SINGLE (world_size = 1) configuration on CPU, and the portable TC_DIST_GLOO TCP baseline now covers localhost multi-rank collectives. TC_DIST_RING remains the v0.5 TB5 transport milestone.
  • Add portable CPU CI coverage — CI builds and tests the CPU-only configuration on every push.
  • New CMake option TC_ENABLE_METAL (defaults: ON on Apple, OFF elsewhere) gates the Apple Metal backend. With TC_ENABLE_METAL=OFF the Metal kernels, MPS fallback, and TensorOps path are excluded; the dispatch ladder collapses to one entry (portable_cpu).
  • Covered ops on the CPU build: tc_init/shutdown, tc_buffer_*, tc_stream_*, tc_gemm (all dtypes + transpose
    • batched + async), attention forward/backward, training kernels, tc_conv2d_forward, tc_quantize_weights, tc_gemv_quantized, tc_gguf_*, tc_dist_* (SINGLE plus portable GLOO TCP), DiLoCo single-rank, dense GLOO, and sparse TOPK GLOO outer steps, CUDA/HIP diagnostics, and diagnostic API (tc_status_string, tc_dtype_name, tc_backend_name).
  • The portable CPU GEMM path now delegates fp32 GEMM and fp16-through-fp32 GEMM to CBLAS when available (Accelerate on macOS, system BLAS on Linux), with the triple-loop implementation retained as the fallback.
  • Portable CPU CI now builds the same installed-SDK shared/static/C++ consumer fixture used by release artifacts, compiles a pkg-config consumer, and runs a Python smoke against the installed shared library (.dylib on macOS, .so on Linux).
  • The Python binding and wheel packaging can load/package a Linux libtensorcore.so for portable CPU builds while preserving the macOS dylib + metallib release contract.
  • Uncovered backend paths (HIP/CUDA execution, dropout tolerance, and non-shipped compression modes) return explicit unsupported statuses so downstream FFI imports can bind the full ABI surface without requiring Metal symbols.

Test surface

  • Add decode and training integration examples — examples/decode_step.c (full synthetic Llama decode step; ~30 ms / 2 layers on M2 Ultra) and examples/training_step.c (RMSnorm + Linear + softmax forward, backward through Linear / RMSnorm, AdamW on weights + gamma; loss decreases 3.98 → 0.25 over 15 steps).
  • Clarify test registration paths — register the new examples as CTest entries (example_decode_step, example_training_step) gated by TC_ENABLE_METAL AND TC_BUILD_TESTS so CPU-only builds skip them.
  • Expand Python GGUF integration coverage — the synthetic GGUF in python/tests/test_basic.py now writes an int64 array (tokenizer.ggml.token_ids) so gguf_meta_array_get_i64 has test coverage. QuantizedMatrix.gemv_async is also exercised against the sync path. Both were previously dead-code-flagged.
  • Test count: 24 / 24 pass (was 22).

CI hardening

  • Harden CI output file writes — scripts/ci_macos_test.sh, scripts/release_smoke.sh, and scripts/create_release_checksums.sh gained require_output_file_path guards that refuse to write to /, /etc, /bin, /usr, /sbin, /System. The checksums script now writes through a mktemp temp file and mvs atomically into place with a named trap handler for cleanup.
  • ICC's shell-hardening audit findings: 3 medium → 0.
  • scripts/check_docs_links.py — auditable check for broken intra-doc Markdown links; wired into the dev Makefile as make docs-check.
  • Release evidence now distinguishes public-core files that were merely not runtime-covered on no-GPU/paravirtual runners (uncovered_files) from files that are actually missing from covered evidence (missing_files).
  • scripts/check_release_evidence.py now enforces that distinction, and CI runs fixture tests for the release-evidence checker.

Dev convenience

  • Makefile with 16 dev targets: build, test, bench, smoke, hello, inspect, decode, train, examples, check-version, check-headers, check-exports, check-python, docs-check, icc-audit, install, wheel. See make help.
  • bench_gemm accepts TC_BENCH_SIZES, TC_BENCH_DTYPES, TC_BENCH_WARMUP, and TC_BENCH_ITERS so public users can run bounded GPU smoke tests or CPU-only reference sweeps without invoking the full 4096³ default pass. CPU-scale throughput now prints as GFLOPS instead of rounding to 0.00 TFLOPS.
  • Pyproject metadata: expanded keywords (13 entries), classifiers (Python 3.9-3.13, C/C++ language tags, Science/Research audience), and [project.urls] with homepage / repo / docs / changelog / roadmap / issues / security URLs.

Documentation overhaul

  • Rebuilt docs/ into a 40+ doc tree, including a CUDA-for-Apple positioning doc (docs/cuda_comparison.md), full C ABI reference (docs/api_reference.md), per-area kernel walkthroughs (gemm, attention, training_kernels, conv2d, quantized, gguf, distributed, dtypes, family_gating), a ground-truth audit produced by ICC, numerics + memory model + observability docs, assembly walkthroughs for inference and training, an FAQ, a glossary, a kernel-extension tutorial, a release runbook, a dev-setup guide, per-directory READMEs in examples/, tests/, bench/, and a per-.metal-file walkthrough in docs/kernels.md.
  • Caught and fixed several docs-vs-code discrepancies along the way: GEMM kernel layout is 4 simdgroups × 128 threads (not 32 × 1024); FlashAttention D=64 tile is Br=Bc=32 (not 64); dgamma from tc_rmsnorm_backward is fp32 (not fp16, with corrected dtype
    • AdamW grad-dtype guidance); tc_backend_name returns lowercase strings, not uppercase enum names.

TensorOps + hardware evidence

  • Align TensorOps device capability gate — the supports_tensorops_m5 gate now mirrors the actual mpp::tensor_ops availability rather than raw GPU family. Older M5 firmware that doesn't expose the encoder reports correctly.
  • Gate TensorOps GEMM to validated tiles, Use SDK26 tensor_ops tile shape, Align tensor_ops GEMM tensors with SDK26 API, Use tensor_ops accumulate mode for SDK26, Probe TensorOps ragged fallback on M5, Add M5 TensorOps runtime smoke wrapper, Disable unvalidated TensorOps attention kernel body — the M5 TensorOps GEMM path is now SDK 26 API-aligned, gated to validated tile shapes, and exercised by a runtime smoke wrapper that is ready for M5 hardware. The TensorOps attention body is intentionally disabled at v0.1 until the validation lands.
  • Add hardware evidence path for TensorOps readiness — release_smoke.sh collects a JSON runtime-evidence artifact (chip, family, TensorOps availability) under a REQUIRE_METAL4_TENSOROPS switch consumed by the self-hosted CI workflow.
  • Add TensorOps build toggle — TC_ENABLE_TENSOROPS now exists as a CMake option. It defaults to ON to preserve SDK 26 compile evidence, but can be set to OFF to force the non-TensorOps Metal path.

Executable end-to-end examples

(See "Test surface" above for the new decode_step.c and training_step.c examples, which are also registered as CTest entries.)

GitHub OSS hygiene

  • SECURITY.md — threat model (GGUF parsing, tensor dimension validation, native lib load), reporting flow, hardening posture.
  • .github/ISSUE_TEMPLATE/{bug_report,feature_request,performance_regression}.md — structured templates that ask for chip / SDK / tc_version() / tc_last_backend() capture up front.
  • .github/PULL_REQUEST_TEMPLATE.md — 11-item checklist citing the numerical guarantees, public-export checks, Python ABI checks, CHANGELOG/ROADMAP/docs updates.
  • CITATION.cff — CFF v1.2.0 metadata for academic / technical citation.

GitHub repo metadata

  • Description set: "CUDA-equivalent tensor-core acceleration for Apple Silicon. C-ABI kernel library wrapping simdgroup_matrix (M1+) and mpp::tensor_ops (M5+): GEMM, FlashAttention, Conv2D, Q4_0/Q8_0 quantized inference, GGUF reader, full transformer training kernels. One binary, M1 → M5."
  • 13 topics added for discoverability: apple-silicon, cuda-alternative, flash-attention, flashattention2, gemm, gguf, llm-inference, metal, mixed-precision, mlx, objective-c-plus-plus, simdgroup-matrix, tensor-ops.

Native SDK release artifact

  • Publish native SDK release artifact — the release workflow now produces a tarball of headers, libraries, and the metallib that downstream consumers can vendor without rebuilding.
  • Added examples/native_sdk_consumer/, a standalone CMake consumer fixture for installed SDKs. Release smoke and native SDK archive verification now build the same shared/static C and C++ consumers that downstream projects can copy directly.
  • Verify native SDK archive consumers, Check static SDK archive consumer, Check C++ SDK archive consumer — CI compiles a tiny consumer against the archived static and shared libraries on every push, in both C and C++ modes, so the archive contract is continuously verified.
  • Publish release checksums + Verify release checksum manifest — release ships a SHA manifest; CI asserts the manifest is consistent with the artifact contents.
  • Harden native SDK archive hygiene — the archive carries no build-tree paths or absolute references; produced via scripts/create_native_sdk_archive.sh and verified by scripts/check_native_sdk_archive.sh.
  • Make SDK CMake exports relocatable — the installed tensorcoreConfig.cmake no longer embeds the build host's prefix; it resolves paths relative to the install root, so the archive works at any prefix on the consumer side.

TensorOps + hardware evidence (extended)

  • Separate packaging evidence from GPU coverage — wheel-packaging and GPU-hardware coverage are now distinct evidence streams so neither masks the other.
  • Harden hardware evidence and Python diagnostics — release smoke emits machine-readable evidence and Python exposes matching diagnostic helpers (status_string, dtype_name, backend_name, last_backend_name, tensorops_gemm_kernel_name).

Distributed and Python ABI

  • Expose distributed primitives in Python — tc_dist_init, tc_allreduce, tc_broadcast, tc_allgather, tc_barrier, tc_dist_world_size, tc_dist_rank plus tc_dist_finalize are now callable from the Python binding.
  • Restrict shared library exports to public ABI — the shared dylib filters its export list down to the symbols declared in include/tensorcore/; CI checks for drift against the public headers on every push (Check public export surface in CI, Check exports against public headers).
  • Public-surface guardrails in CI — Check Python FFI surface, Check Python ABI layout, Check Python constants against public enums, Check public headers. The Python binding can no longer drift from the C ABI undetected.

Release workflow

  • Use release smoke for release workflow + Handle paravirtual GPUs in release smoke — release.yml now drives scripts/release_smoke.sh, which copes with paravirtualized GPUs that show up under macOS virtualization.

v0.1.22 — Version guards across pyproject / CMake / header

  • Added scripts/check_version_consistency.sh to fail CI if pyproject.toml, CMakeLists.txt::project(VERSION ...), and the TENSORCORE_VERSION_{MAJOR,MINOR,PATCH} triple in include/tensorcore/tensorcore.h disagree.
  • The CI workflow runs the check as its first step on every push.

v0.1.21 — Release wheel workflow + macOS hardware-aware CI

  • Add release wheel workflow — .github/workflows/release.yml builds the tensorcore_apple-*.whl on the v* tag, runs the version check, hardware-aware tests, the install smoke, and publishes the wheel as a GitHub release asset.
  • Upload release wheel artifacts — wheels also upload as build artifacts on every successful release run.
  • Update checkout action for Node 24 — bumps actions/checkout to v6 for Node 24 compatibility.
  • Use a venv for Python CI smoke — the Python smoke script now sets up an isolated venv so system Python state can't leak in.
  • Make macOS CI hardware-aware — scripts/ci_macos_test.sh skips tests that require a real GPU when the runner only exposes the paravirtual one, so CI on macos-14 / macos-15 runners is predictable.

v0.1.20 — Public integration oracle

  • Add public integration oracle — a CI-side check that compiles a minimal tc_gemm consumer against the installed CMake package, runs it, and asserts the expected backend dispatch. The oracle is the contract for "the installed package is usable from another project."

v0.1.19 — TensorOps selector extracted for coverage

  • Extract TensorOps selector for coverage — the M5 mpp::tensor_ops path selection moves to lib/tensorops/tensorops_select.{c,h} so it can be unit-tested independently of a Metal 4 runtime.
  • tests/test_tensorops_select.c covers the selector's dtype × accum matrix without needing M5 hardware.

v0.1.18 — Runtime evidence for public integration

  • Expand runtime evidence for public integration — release_smoke.sh emits structured runtime evidence (backend chosen, family detected, TensorOps presence) so downstream integrators can see what their installed binary does on real hardware, not just what it claims.

v0.1.17 — Conv2D backward in Python + runtime evidence

  • Expose Conv2D backward and runtime evidence — tc_conv2d_backward_input and tc_conv2d_backward_weight are wrapped in the Python binding; runtime-evidence helpers (conv2d_backward_input_scratch_bytes) ship alongside them.

v0.1.16 — Python buffer layout hardening

  • Harden Python buffer layout handling — buffer_write / buffer_read validate dtype, shape, and contiguity against the buffer they're targeting; mismatched layouts raise instead of silently corrupting data.

v0.1.15 — Wheel native artifact tags validated

  • Validate wheel native artifact tags — the release pipeline asserts the wheel's platform tag matches the dylib it carries (no macosx_15_0_arm64.whl shipping a dylib built for macosx_11_0).

v0.1.14 — Python native library loading hardened

  • Harden Python native library loading — the binding's _find_lib searches in a predictable order (TENSORCORE_LIB env → next to the package → standard install prefixes → build tree) and produces a clear TensorcoreError if every candidate fails, with the searched paths in the message.

v0.1.13 — Conv2D forward in Python

  • Expose Conv2D forward through Python — tc_conv2d_forward is wrapped in the Python binding, with helpers (conv2d_output_shape, conv2d_scratch_bytes) that match the C-side scratch-sizing math.

v0.1.12 — Batched GEMM in Python

  • Expose batched GEMM through Python — tc_gemm_batched is wrapped; the binding takes a single tc_gemm_batched_desc and exposes the stride parameters cleanly.

v0.1.11 — Attention in Python

  • Expose attention through Python bindings — tc_attention_forward, tc_attention_forward_async, and tc_attention_backward cross the Python boundary. The descriptor (including kv_heads, window_size, alibi_slopes) is exposed; alibi_slopes accepts a NumPy fp32 array.

v0.1.10 — Async attention dispatch hardened

  • Harden async attention dispatch — the _async attention variant shares the same pending command-buffer pattern as tc_gemm_async so multiple attention dispatches can pipeline through one stream.

v0.1.9 — Buffer validation + wheel release guards

  • Bump version to 0.1.9 — first wheel-publishable release.
  • Validate buffers before GPU dispatch — tc_gemm, tc_attention_forward, and friends now validate that the supplied tc_buffer* objects come from the same context and have sufficient size for the requested shape, returning TC_ERR_INVALID_ARG / TC_ERR_INVALID_SHAPE instead of trusting the caller.
  • Harden Python wheel release checks — wheel release adds version / artifact assertions before upload.

v0.1.8 — Native runtime packaged in wheels

  • Package native runtime in Python wheels — the wheel now ships libtensorcore.dylib and tensorcore.metallib inside the package, so pip install tensorcore-apple is a complete install on Apple Silicon. TENSORCORE_LIB / TC_METALLIB env overrides still take precedence.

v0.1.7 — pkg-config support + owned Python wrappers

  • Add pkg-config install support — tensorcore.pc ships in lib/pkgconfig/. Direct compiler invocation works: cc main.c $(pkg-config --cflags --libs tensorcore).
  • Check pkg-config metadata in CI — CI validates the generated tensorcore.pc against the install prefix on every push.
  • Add owned Python context and buffer wrappers — tc.Context, tc.Buffer, and tc.Stream are context-manager-friendly object wrappers that own the underlying handle and release on exit.
  • Add owned Python GGUF wrappers — tc.GgufFile, tc.LoadedModel, tc.LoadedTensor, tc.QuantizedMatrix mirror the pattern for the GGUF surface.
  • Add NumPy helpers to Python buffers — Buffer.from_numpy, Buffer.to_numpy, Context.buffer_from, Context.buffer_zeros.
  • Fix context and Python ownership lifetimes — wrappers correctly retain the context as long as any dependent handle is alive, so GC ordering can't pull the rug out from under a live buffer.
  • Support installed static consumers — out-of-tree builds linking against tensorcore::tensorcore (the static library) resolve the Metal / Foundation / Accelerate frameworks correctly via the installed CMake package config.

v0.1.6 — GGUF loading, Q4 v2, Q8 quantization, installable package

This checkpoint turns tensorcore from a build-tree kernel library into something downstream projects can start consuming directly.

GGUF reader and bulk tensor loading

  • Added include/tensorcore/gguf.h and lib/io/gguf.c: a memory-mapped GGUF v3 reader with metadata lookup, tensor enumeration, and tensor-to-buffer copy.
  • Added tc_gguf_load_supported_tensors: bulk-copies every supported GGUF tensor into owned tc_buffer objects and reports skipped unsupported encodings.
  • Added GGUF numeric and array metadata helpers for model config and tokenizer fields: strings, integers, floats, string arrays, integer arrays, and float arrays.
  • Added tc_gguf_get_llama_config / Python gguf_get_llama_config to extract common LLaMA-family dimensions, head counts, RoPE settings, RMSNorm epsilon, and vocab size in one call.
  • Added GGUF quantized matrix descriptor helpers in C and Python so runtimes can safely map GGUF [K, N] Q4_0/Q8_0 tensors to tc_gemv_quantized(..., N, K).
  • Added tests/test_gguf.c: synthetic GGUF round-trip, metadata validation, direct tensor copy, bulk load, unsupported-tensor skip count, and Q4 GEMV from a GGUF-backed tensor.
  • Added examples/gguf_inspect.c: CLI inspection plus explicit --load-supported GPU copy mode.

Quantized inference

  • Added kernels/metal/gemm_quantized_v2.metal: faster Q4_0 GEMV path used by default; TC_Q4_USE_V1=1 keeps the older kernel available for comparison.
  • Fixed Q4_0 packing to GGML/GGUF layout: low nibble is weight i, high nibble is weight i+16.
  • Added GPU Q8_0 quantization via tc_quantize_q8_0; public tc_quantize_weights(..., TC_QUANT_Q8_0, ...) now works.
  • tests/test_quantized.c now covers Q4_0 sync/async, tail N, Q8_0 GPU quantize+GEMV, and invalid quant enum sizing.

Streams and installability

  • Stream-backed async ops now share a pending command buffer so batched inference calls avoid per-op command-buffer round trips.
  • Added installable CMake package export: tensorcore::tensorcore and tensorcore::tensorcore_shared.
  • Added installed lib/pkgconfig/tensorcore.pc for Makefile/direct-compiler consumers.
  • Installed builds now ship lib/tensorcore.metallib, and the runtime finds it next to the loaded dylib without requiring TC_METALLIB.
  • Added docs/integrating_tensorcore.md with the CMake, C ABI, Python, and GGUF bulk-load integration paths.

Python binding

  • Expanded python/tensorcore/__init__.py beyond GEMM: streams, async GEMM, RMSNorm/LayerNorm/RoPE/SwiGLU/softmax/AdamW/fused RMSNorm+GEMV wrappers, Q4/Q8 quantized helpers, GGUF metadata/tensor access, single tensor copy, matrix descriptors, and bulk loaded-model handles.
  • Added pyproject.toml so downstream projects can install the Python binding with pip while using a CMake-installed libtensorcore.dylib.
  • python/tests/test_basic.py now validates GEMM, async GEMM, training wrappers, Q4_0, Q8_0, GGUF copy, and GGUF bulk load.
  • CTest registers python_basic when Python + NumPy are available.

Verification

  • ctest --test-dir build --output-on-failure: 17/17 pass on Apple M2 Ultra.
  • Installed package smoke verified with /private/tmp/tensorcore-install.
  • Out-of-tree CMake consumer verified via find_package(tensorcore CONFIG).

v0.1.5 — Modern attention, multi-process distributed, inference bench, Python

Closes the rest: sliding-window + ALiBi attention (modern LLM features), real multi-process ring all-reduce via fork(), a synthetic 7B Q4_0 inference latency bench with concrete tok/s, and a Python ctypes binding so the library is usable from any numpy script.

Sliding-window + ALiBi attention

  • New tc_attention_desc fields: window_size (Mistral-style local attention, 0 = full attention) and alibi_slopes (per-head ALiBi linear bias, host fp32 array).
  • kernels/metal/flash_attention.metal: applies window + ALiBi in the score-modification step under new function constants g_use_window / g_use_alibi. Causal + window + ALiBi can be combined.
  • tests/test_attention_correctness.c: added sliding-window case (W=16 with Sq=Sk=64). Validated at 1.6e-3 RMS-scaled vs fp64 reference.

Multi-process ring all-reduce (real fork)

  • tests/test_distributed_ring_fork.c: same ring algorithm as the threads test, but each rank is a fork()ed child. Communicates via socketpairs the parent set up before fork. Validated bit-exact for 4 ranks × 1024 fp32 elements.
  • This is the same code path the multi-Mac TB5/RDMA backend will use — only the transport layer changes.

Synthetic 7B Q4_0 inference bench

  • bench/bench_inference_7b.c: allocates Q4_0 weights matching a 7B llama (32 layers × hidden=4096 × mlp=11008 = 3.4 GB), times the GEMV stack per decode step. Excludes attention/softmax/RoPE/RMSnorm.
  • Initial result (sync-per-call): 6.5 tok/s @ 22 GB/s.
  • After async batched dispatch: 13.7 tok/s @ 46.5 GB/s.
  • Reference: llama.cpp on M2 Ultra reports ~55-65 tok/s. We're at 24% of that today; the gap is in Q4_0 kernel design (1 sg per output cell with no inter-block pipelining vs llama.cpp's hand-tuned 4-output-per-sg).
  • Added tc_gemv_quantized_async for stream-batched dispatch.

Python ctypes binding

  • python/tensorcore/__init__.py: minimal ctypes wrapper for tc_init, tc_buffer_*, tc_gemm with TCDeviceInfo / TCGemmDesc Python structs. numpy interop via buffer_write / buffer_read.
  • python/tests/test_basic.py: end-to-end fp16 GEMM 256³ vs numpy reference, validates the binding is functional.
  • CMake now builds both libtensorcore.a and libtensorcore.dylib; the .dylib is what ctypes loads.

Eshkol integration in BOTH repos

  • ~/Desktop/eshkol-platform/ was the wrong repo earlier; the user clarified the main branch is ~/Desktop/eshkol/. Both now have the same integration (drop the bridge file, one-line call site, build clean). Separate commits in each repo. eshkol-static builds green in both.

Test count: 15/15 pass on Apple M2 Ultra

Added test_distributed_ring_fork (thread → process), sliding-window case in test_attention_correctness, test_quantized + test_fused_norm_gemv from prior. Total ctest time ~3s.

v0.1.4 — Quantized inference + GQA + fused norm + REAL Eshkol integration

The big push: the LLM-inference kernels (Q4_0/Q8_0 weight-only matmul), GQA attention validation, fused RMSnorm+GEMV for the inference hot path, and the long-deferred actual integration with eshkol-platform.

Q4_0 / Q8_0 quantized matmul (LLM inference)

  • kernels/metal/gemm_quantized.metal: ggml-style block quantization. Q4_0 = 32 weights/block × (fp16 scale + 16 packed nibbles) = 4.5 bits/weight. Q8_0 = 32 weights/block × (fp16 scale + 32 int8) = 8.5 bits/weight.
  • tc_q4_0_gemv_f16, tc_q8_0_gemv_f16: dequantize on the fly, multiply against fp16 activation, output fp16. One simdgroup per output cell, cooperative simd_sum reduction.
  • tc_quantize_q4_0: GPU-side quantization kernel (rounds fp16 weights into Q4_0 blocks with per-block scale).
  • Public API in include/tensorcore/quantized.h: tc_quantize_weights, tc_gemv_quantized, tc_quantized_size.
  • tests/test_quantized.c: validated bit-exact against CPU dequant ref (rms_scaled=2.0e-4); storage = exact 4.50 bits/weight at K=256.

GQA / MQA attention validation

  • tests/test_attention_correctness.c: added 3 GQA cases (MQA with 1 KV head, GQA with H/2 KV heads, GQA H=8 KV=2 with D=128). All pass at <1% RMS-scaled error vs fp64 reference that implements the same H→KV_H head-grouping the kernel does.

Fused RMSnorm + GEMV

  • kernels/metal/fused_norm_gemv.metal: one-pass Y = RMSnorm(X, gamma) @ W. Eliminates the round-trip of the normalized intermediate — the dominant cost at inference batch sizes (M≤4). Two-pass intra-threadgroup: pass 1 computes rstd via simd reductions; pass 2 reapplies normalization inline as part of the matmul accumulation.
  • tc_fused_rmsnorm_gemv public API.
  • tests/test_fused_norm_gemv.c: validated by comparing fused output to tc_rmsnorm_forward + tc_gemm separate path, rms_scaled<5e-3.

REAL eshkol-platform integration (the long-promised one)

  • The eshkol/bridge/tensorcore_codegen.cpp shim was previously only compile-tested standalone. Now it's actually dropped into eshkol-platform/lib/backend/, glob-included by their CMakeLists, and called from the codegen-context initialization path.
  • Activation is opt-in via ESHKOL_ENABLE_TENSORCORE=1. When set, the 14 tc_* C ABI functions are declared as ExternalLinkage in the Eshkol LLVM module at codegen-context init time.
  • eshkol-platform builds 100% clean with the bridge in place (eshkol-static target green, REPL functional in both modes: (+ 1 2) → 3 identically with and without the env var).
  • Eshkol-platform changes: 1 file added + 1 line change at the call site. Fully reversible — set env=0 or remove the file and the build is back.

Test count: 14/14 pass on Apple M2 Ultra

Added test_quantized, test_fused_norm_gemv (12→14). All prior tests still pass. ctest --output-on-failure completes in ~3s.

v0.1.3 — Universal-dtype GEMM + multi-batch Conv + macOS 26 SDK gating

Closes every remaining hardware-gated path from v0.1.2 by adding software fallbacks that work on every M-series chip today. bf16 and i8 GEMM no longer require M3+/M4+ — they validate on this M2.

Software bf16 + i8 GEMM (every M-series, today)

  • lib/fallback/mps_gemm.mm: added bf16_via_fp32 and i8_via_fp32. The bf16 path bit-casts bf16↔fp32 (bf16 = high 16 bits of fp32) and routes through tc_gemm fp32. The i8 path is exact (fp32 has 24-bit mantissa, more than enough for int8·int8 sums up to K=2^16).
  • tests/test_gemm_bf16.c: was skipping on Apple<9; now runs and passes on M2 Ultra at all 4 shapes. RMS-scaled error ~2.7e-3 vs fp64 reference.
  • tests/test_gemm_i8.c: was skipping on Apple<10; now runs and passes. Bit-exact (0 errors across 65K cells at 256³).

Multi-batch Conv2D backward input

  • lib/ops/conv.mm tc_conv2d_backward_input: per-batch GEMM with MTLBuffer offset binding, mirrors the dW pattern. Validated by test_conv2d which now uses N=1 but the code path scales to N>1.

128×128 async tile (env-gated experimental)

  • kernels/metal/gemm_async_128.metal: written + compiled. Currently regresses perf on M2 (~10 vs ~19 TFLOPS at 4096³) due to 16-frag/sg register pressure. Opt-in via TC_USE_ASYNC_128=1 for benchmarking; expected to win on M3+/M4 with more registers per simdgroup.

macOS 26 forward-compat

  • CMake auto-detects SDK version and gates gemm_async.metal / gemm_async_128.metal out of the build when SDK >= 26.0 (Xcode 17+ rejects the __asm("air.simdgroup_async_copy_2d.…") form per the AGX ISA research). Build succeeds on either SDK; dispatch logic runtime-probes the metallib symbol and silently falls back to the sync vec4 path when async kernels aren't present.

Measured perf on M2 Ultra (Apple8, ~27 TFLOPS theoretical)

Workload TFLOPS % peak Notes
fp16 GEMM 4096³ async 19.30 72% async_copy via private AIR intrinsics
fp32 GEMM 4096³ 2.43 60% bit-exact vs Accelerate
bf16 GEMM (SW path) matches fp32 minus quantization n/a new in v0.1.3
i8 GEMM (SW path) bit-exact int32 n/a new in v0.1.3

Test count: 12/12 pass on Apple M2 Ultra

All tests run end-to-end on this hardware now. Nothing "skips cleanly because silicon lacks feature" anymore.

v0.1.2 — Async DMA + real distributed + Conv tests

Closes everything I deferred in v0.1.1. No more "this is gated by hardware" — kernels validated, paths exercised end-to-end.

Major: simdgroup_async_copy in GEMM (the perf prize)

  • kernels/metal/metal_simdgroup_event.h: shim header declaring the private AIR intrinsics (air.simdgroup_async_copy_2d.p3i8.p1i8, air.wait_simdgroup_events) reverse-engineered by the Philip Turner / MFA effort. C++ wrapper class tc::simdgroup_event mirroring the MFA API.
  • kernels/metal/gemm_async.metal: GEMM that issues async DMAs from sgid==0, waits via simdgroup_event::wait(2, ev), barrier-publishes to peer simdgroups, computes. Single-buffered (MFA pattern, not double-buffer).
  • Opt-in via TC_USE_ASYNC=1. Measured on M2 Ultra:
Shape sync (vec4) async delta
4096³ fp16 17.65 TFLOPS 18.99 TFLOPS +7.6%
2048³ fp16 10.05 11.86 +18%
1024³ fp16 3.12 4.38 +40%
  • Compatibility note in the shim header: macOS 26+ / Xcode 17+ rejects the __asm("air.…") form. v0.2 will ship the AIR-IR fallback the way MFA does.

Major: real ring all-reduce

  • lib/distributed/ring_local.mm: full Rabenseifner ring (reduce-scatter + all-gather) over socketpair(AF_UNIX, SOCK_STREAM). The transport-swap to multi-Mac TB5 (or RDMA verbs via librdma.tbd) is a single function point.
  • tc_dist_ring_pair_make: build N socketpair-connected ring edges.
  • tc_dist_ring_local_allreduce_ex: bandwidth-optimal algorithm, fp32-sum + fp16-sum implemented; per-rank traffic is 2(N-1)/N · |B|.
  • tests/test_distributed_ring.c: WORLD=4 threads, N_ELEMS=1024 fp32 sum. Validated bit-exact against single-process sum (max_abs_err=0).

Major: Conv2D correctness + multi-batch dW

  • tests/test_conv2d.c: forward validated vs fp64 CPU reference (rms_scaled=3.97e-04). Backward input + weight kernels both dispatch and write nonzero results.
  • lib/ops/conv.mm tc_conv2d_backward_weight: now loops over batches with beta=1 accumulation on subsequent iterations. Replaces the v0.1.1 stub that silently computed only batch 0.

Research deliverables

  • Two deep-dive research reports informed the work above:
    • simdgroup_async_copy API — confirmed exists, found MFA's pattern, debunked my prior "Metal has no async DMA" claim.
    • Distributed Metal landscape — JACCL/TB5/RDMA via librdma.tbd, MLX ring source patterns, IOSurface+MTLSharedEvent for cross-process GPU buffers.

Test count: 12/12 pass on Apple M2 Ultra

  • test_device, test_gemm_f32, test_gemm_f16, test_gemm_bf16, test_gemm_i8,
  • test_attention_correctness, test_attention_backward, test_training_kernels,
  • test_transformer_block, test_e2e_training, test_conv2d, test_distributed_ring

Cumulative GEMM perf trajectory on M2 Ultra (~27 TFLOPS theoretical peak)

Version fp16 4096³ % peak What changed
v0.1.0 initial 13.75 51% basic simdgroup_matrix, scalar loads
v0.1.0 + vec4 16.46 61% vec4 cooperative loads
v0.1.0 + BK=32 17.59 65% larger K-block per iteration
v0.1.2 + async_copy 18.99 70% MFA-style async DMA

Still chasing MLX (~21 TFLOPS, ~78% peak); the remaining gap is in epilogue scheduling + register-pressure-aware 128×128 tile (v0.2).

v0.1.1 — Training-complete

Adds the rest of the training stack on top of v0.1.0's kernel substrate.

New kernels

  • flash_attention_backward_d128.metal: FlashAttention backward at head_dim=128 (Br=Bc=16, fits 32 KB TG mem). dQ + split dK/dV kernels. Validated <1% RMS-scaled error vs fp64 reference.
  • gemm_simdgroup.metal: added tc_gemm_f16_f32_batched — single-kernel batched fp16 GEMM with per-batch strides. Replaces the per-batch host loop for fp16 alpha=1/beta=0 cases.
  • conv2d_backward.metal: tc_col2im_atomic_f32 (scatter-add via fp32 atomics) and tc_col2im_finalize_f16 (fp32→fp16).

New host APIs

  • tc_attention_backward now handles D=128 in addition to D=64; same tc_attention_desc interface, head_dim picks the kernel variant.
  • tc_gemm_batched fast path on fp16: single dispatch with MTLSize(gx, gy, batch). Falls back to per-batch loop for other dtype/transpose configs.
  • tc_conv2d_backward_input (col2im scatter-add path), tc_conv2d_backward_weight (im2col + GEMM with transpose_b).
  • Bench-driven autotune wired at tc_init: TC_AUTOTUNE=1 triggers a one-time probe that caches the per-device tile config to ~/.tensorcore/autotune_<device>.json and reloads on subsequent runs.

Eshkol integration validated

  • eshkol/bridge/tensorcore_codegen.cpp now ships with compile evidence: object file produced cleanly against eshkol-platform/inc/eshkol/backend/codegen_context.h + Homebrew LLVM. nm confirms _eshkol_register_tensorcore_builtins is an exported global symbol. See eshkol/bridge/COMPILE-EVIDENCE.txt.

New tests (10/10 pass on Apple M2 Ultra)

  • test_e2e_training: real multi-step training loop. MLP memorizes a random target via 100 AdamW steps. Loss 8.37e-2 → 2.60e-5 (100% reduction). Exercises GEMM forward, SwiGLU, GEMM with transpose_a and transpose_b for backward, AdamW fp32-master/fp16-grad update path.
  • test_attention_backward extended with D=128 case.

Known not-yet-shipped (deferred to v0.2)

  • simdgroup_async_copy MFA-style pattern adoption in GEMM. Compile-time gate (TC_HAVE_ASYNC_COPY) is in but the kernel still uses vec4 cooperative loads. Avoiding this in v0.1 because Metal lacks an explicit async DMA primitive (verified via dougallj/applegpu research) and the prior double-buffer attempt regressed perf. Real path requires M3+ hardware to validate the explicit async copy.
  • bf16 / int8 perf validation (M2 Ultra silicon doesn't expose those simdgroup_matrix variants; kernels compile and dispatch-skip cleanly).
  • Multi-batch Conv2D forward and dW accumulation (single-batch only on this path).
  • Real Thunderbolt-5 ring + JACCL distributed backend (single-host emulation is live; multi-Mac is a phase v0.5 hardware-validation milestone).

v0.1.0 — Foundation

Kernels (Metal)

  • gemm_simdgroup.metal: 64×64 GEMM, BK=32, vec4 cooperative loads, fp16/bf16/fp32/i8 with fp32 accumulators
  • gemm_simdgroup_128.metal: 128×128 large-tile variant (opt-in via TC_USE_128_TILE=1)
  • flash_attention.metal: fused FA-2 forward, D=64
  • flash_attention_d128.metal: fused FA-2 forward, D=128
  • flash_attention_backward.metal: split-kernel dQ + dK/dV backward (D=64)
  • training_kernels.metal: RMSnorm fwd+bwd, LayerNorm fwd+bwd, RoPE fwd, SwiGLU fwd+bwd, softmax fwd+bwd, fused AdamW step
  • conv2d.metal: im2col + bias-add (forward, via tc_gemm)
  • tensorops_gemm.metal: Metal 4 mpp::tensor_ops::matmul2d path (SDK 26+, M5 Neural Accelerator)
  • tensorops_flash_attention.metal: Metal 4 FlashAttention skeleton (SDK 26+, validation pending M5)

Public C ABI

  • Lifecycle: tc_init, tc_shutdown, tc_device_info_get, tc_version
  • Buffers: tc_buffer_alloc, tc_buffer_free, tc_buffer_map, tc_buffer_size
  • Streams: tc_stream_create, tc_stream_destroy, tc_stream_sync
  • GEMM: tc_gemm, tc_gemm_async, tc_gemm_batched
  • Attention: tc_attention_forward, tc_attention_forward_async, tc_attention_backward
  • Training: tc_rmsnorm_forward/_backward, tc_layernorm_forward/_backward, tc_rope_forward, tc_swiglu_forward/_backward, tc_softmax_forward/_backward, tc_adamw_step
  • Conv: tc_conv2d_forward
  • Distributed: tc_dist_init, tc_dist_finalize, tc_allreduce, tc_broadcast, tc_allgather, tc_barrier (single-host backend live; ring TB5 + Gloo gated for v0.5)
  • Diagnostics: tc_last_backend, tc_backend_name, tc_status_string, tc_dtype_name

Runtime

  • lib/core/device.mm: Apple GPU family detect (Apple7..Apple11) + unified-memory probe
  • lib/core/pipeline_cache.mm: thread-safe MTLComputePipelineState cache, function-constant specialization
  • lib/core/buffer_pool.mm: power-of-2 bucketed MTLBuffer pool (LIFO recycle, 8/bucket cap)
  • lib/core/autotune.cpp: family-keyed tile selection + cache load/save
  • lib/tensorops/tensorops_m5.mm: Metal 4 host dispatch (SDK-gated)
  • lib/distributed/distributed.mm: single-host backend; TB5 ring + Gloo stubs

Fallbacks

  • lib/fallback/mps_gemm.mm: MPSMatrixMultiplication path
  • lib/fallback/accelerate_gemm.c: CPU cblas_sgemm (AMX on M1-M3, SME on M4+)

Tests (9 total, 100% passing on M2 Ultra)

  • test_device: smoke + family detect
  • test_gemm_f32: bit-exact vs Accelerate (max_abs=0 across all shapes)
  • test_gemm_f16: RMS-scaled error vs Accelerate, <1.5e-2 across 64..512
  • test_gemm_bf16: kernel skip-clean on Apple<9 (no runtime exercise on M2)
  • test_gemm_i8: kernel skip-clean on Apple<10
  • test_attention_correctness: D=64 and D=128 vs fp64 reference, <2e-2 RMS-scaled
  • test_attention_backward: dQ/dK/dV all <1% RMS-scaled vs fp64 analytic gradient
  • test_training_kernels: 6/6 kernels (RMSnorm/LayerNorm/SwiGLU/softmax/RoPE/AdamW)
  • test_transformer_block: full forward through every kernel + AdamW step

Measured perf (Apple M2 Ultra, family Apple8, ~27 TFLOPS fp16 peak)

Workload TFLOPS % of peak
GEMM fp16 4096³ 17.59 ~65%
GEMM fp32 4096³ 2.38 ~60% (bit-exact)
FA fwd fp16 D=64 S=4096 6.72 —

Eshkol integration

  • eshkol/tensorcore.esk: Scheme-level bindings (tc-init, tc-gemm-fp16, etc.)
  • eshkol/hello_tensorcore.esk: minimal example
  • eshkol/bridge/tensorcore_codegen.cpp: drop-in for eshkol-platform/lib/backend/ — declares 14 tc_* ExternalLinkage LLVM symbols, mirrors builtin_declarations.cpp pattern
  • eshkol/bridge/INTEGRATION.md: 4-step recipe

Build

  • CMake 3.20+, macOS 12.0+ (Apple7+ runtime check). C11/C++17.
  • SDK detection auto-includes Metal 4 sources when SDK >= 26.0; skipped cleanly on older SDKs (today: macOS 15.1 + Xcode 16.2 + SDK 15.2).
  • compile_metallib.cmake helper: .metal → .air → default.metallib precompile (qgt-style, no runtime compile overhead).

Known limitations (documented in ROADMAP.md)

  • v0.1 bf16/i8 paths unexercised at runtime (M2 lacks the silicon).
  • v0.1 conv2d covers forward only and processes batches serially.
  • v0.1 distributed: only single-host backend live; multi-Mac TB5 ring lands v0.5.
  • v0.1 attention backward: D=64 only.
  • v0.1 autotune: family-keyed static table; bench-driven sweep + cache persistence are wired but not yet self-tuning at init.
  • Metal 4 mpp::tensor_ops attention kernel has placeholder softmax step pending M5 hardware validation.