rlx-metal 0.2.14

Metal backend for RLX — Apple Silicon GPU via Metal Performance Shaders + custom MSL kernels
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168

// RLX — versatile ML compiler + runtime.
// Copyright (C) 2026 Eugene Hauptmann, Nataliya Kosmyna.
// SPDX-License-Identifier: MIT OR Apache-2.0

// `objc` crate's `class!` / `msg_send!` macros expand to
// `cfg(feature = "cargo-clippy")` checks that aren't recognized by
// modern rustc. The warnings are third-party noise (~78 across this
// crate); they say nothing about our code. Silence at the crate root.
#![allow(unexpected_cfgs)]

//! RLX Metal backend — Apple Silicon GPU execution.
//!
//! Compiles RLX IR graphs to Metal compute pipelines + MPS matrix kernels.
//!
//! Architecture mirrors rlx-cpu:
//! - `device` — Metal device discovery and properties
//! - `arena`  — GPU buffer allocation from memory plan
//! - `blas`   — MPS matrix multiplication (analog of cblas_sgemm)
//! - `kernels`— custom MSL compute shaders (analog of NEON kernels)
//! - `thunk`  — pre-compiled command buffer with arena offsets
//! - `backend`— ExecutableGraph implementation
//!
//! Apple Silicon advantages:
//! - Unified memory: zero-copy CPU↔GPU
//! - 16-core GPU on M4 Pro: ~1.4 TFLOP/s peak
//! - 273 GB/s memory bandwidth (vs 120 on CPU)
//! - MPSMatrixMultiplication uses dedicated matmul hardware

#[cfg(rlx_metal_host)]
pub mod device;

#[cfg(rlx_metal_host)]
pub mod arena;

#[cfg(rlx_metal_host)]
pub mod blas;

#[cfg(rlx_metal_host)]
pub mod mps_blas;

#[cfg(rlx_metal_host)]
pub mod mps_graph;

#[cfg(rlx_metal_host)]
pub mod mps_graph_hybrid;
#[cfg(rlx_metal_host)]
pub mod mps_graph_lower;

#[cfg(rlx_metal_host)]
pub mod mps_gelu;

/// Hand-rolled Metal bindings, replacing the `metal` crate. Public so downstream
/// `MetalGpuKernel` authors can name the encoder/buffer types they are handed:
/// `use rlx_metal::mtl as metal;` keeps existing custom kernels source-compatible.
#[cfg(rlx_metal_host)]
pub mod mtl;

#[cfg(rlx_metal_host)]
pub mod icb;

#[cfg(rlx_metal_host)]
pub mod kernels;

pub mod vmath;

#[cfg(rlx_metal_host)]
pub mod fft_dispatch;

/// CPU host-fallback for the core Riemannian / SPD-manifold ops (BiMap /
/// ReEig / LogEig / SpdBatchNorm / SpdKarcherMean + backwards). No MSL
/// eigen kernel; they run `rlx_cpu::spd` (F64) against the unified-memory
/// arena between GPU segments, like `Op::Fft`. See `crate::spd`.
#[cfg(rlx_metal_host)]
pub mod spd;

#[cfg(rlx_metal_host)]
pub mod hc_sinkhorn_gate;
#[cfg(rlx_metal_host)]
pub mod llada2_gate;
#[cfg(rlx_metal_host)]
pub mod ms_deform_attn;

#[cfg(rlx_metal_host)]
pub mod config;
#[cfg(rlx_metal_host)]
pub use config::{
    MetalRuntimeConfig, install_runtime_config, reload_runtime_config, runtime_config,
};

#[cfg(rlx_metal_host)]
pub mod cost;

#[cfg(rlx_metal_host)]
pub mod calibrate;

#[cfg(rlx_metal_host)]
pub mod thunk;

#[cfg(rlx_metal_host)]
pub mod backend;

#[cfg(rlx_metal_host)]
pub mod attention_bwd_gpu;

#[cfg(rlx_metal_host)]
pub mod thunk_profile;

#[cfg(rlx_metal_host)]
pub mod mps_profile;

#[cfg(all(feature = "native-splat", rlx_metal_host))]
pub mod splat_adam;
#[cfg(all(feature = "native-splat", rlx_metal_host))]
pub mod splat_native;
#[cfg(all(feature = "native-splat", rlx_metal_host))]
pub mod splat_training;
#[cfg(all(feature = "native-splat", rlx_metal_host))]
pub mod splat_training_pipeline;

/// Double-single (2× f32 ≈ f64) reductions — near-f64 precision on Metal, which
/// has no native f64. Compiled with precise math (fast-math breaks EFT).
#[cfg(rlx_metal_host)]
pub mod double_single;
#[cfg(rlx_metal_host)]
pub mod pipeline_cache;

#[cfg(rlx_metal_host)]
pub mod onnx_qmatmul;

#[cfg(rlx_metal_host)]
pub mod async_copy;

#[cfg(rlx_metal_host)]
pub mod op_registry;

#[cfg(rlx_metal_host)]
pub mod prefill_stats;

#[cfg(rlx_metal_host)]
pub mod collective;

/// Legalization op claim — always available (no Metal device required).
pub mod supported_ops;
pub use supported_ops::SUPPORTED_OPS;

/// PLAN: Schedule splitting for the Metal MPSGraph path. Splits the
/// schedule at attention boundaries so the broken slice-of-computed
/// MPSGraph attention pattern is replaced by the parity-correct
/// thunk path; everything else still gets the MPSGraph dispatch-
/// overhead reduction. Scaffolding only today (data model +
/// segmenter + 3 unit tests); executor wiring + per-segment plan
/// compilation is the next chunk.
pub mod segmented;

/// Typed kernel planning, precision legality, token bucketing, and two-stage
/// build/launch configuration for Metal kernel families.
pub mod kernel_plan;

/// Whether a usable Metal device is present. `rlx-metal` is a Metal-only
/// dependency (its consumers gate it to `cfg(all(target_vendor = "apple",
/// not(target_os = "watchos")))` — every Apple platform with Metal: macOS,
/// iOS, tvOS, visionOS), so this is never a non-Apple `false` stub — callers
/// on other platforms (and watchOS) report Metal availability via the
/// runtime's own device-feature check, not this crate.
#[cfg(rlx_metal_host)]
pub fn is_available() -> bool {
    device::has_metal_device()
}