Skip to main content

rlx_ir/
lib.rs

1// RLX — versatile ML compiler + runtime.
2// Copyright (C) 2026 Eugene Hauptmann, Nataliya Kosmyna.
3// SPDX-License-Identifier: MIT OR Apache-2.0
4
5//! RLX Tensor IR — the intermediate representation for the RLX ML compiler.
6//!
7//! This IR is:
8//! - **Standalone**: no runtime, no backend, no framework coupling
9//! - **Serializable**: graphs can be saved/loaded for AOT compilation
10//! - **Optimizable**: designed for pattern-matching fusion and buffer planning
11//!
12//! The compiler pipeline has three named levels:
13//!
14//! - **HIR** ([`hir`]) — block-oriented IR for model builders (`Linear`,
15//!   `SwiGLU`, `ResidualRmsNorm`, …).
16//! - **MIR** ([`mir`]) — fused tensor DAG; input to [`rlx_opt`].
17//! - **LIR** ([`lir`]) — optimized MIR + arena buffer plan for backends.
18//!
19//! [`Graph`] is the primary DX surface. Use [`Graph::define`] for
20//! fusion-first HIR builders, or [`Graph::new`] / [`GraphModule::mir`]
21//! for primitive MIR. [`GraphModule`] tracks pipeline stage (HIR/MIR/LIR).
22//!
23//! - [`Graph`]: a DAG of tensor operations (like XLA's HloModule)
24//! - [`Node`]: a single operation with typed inputs/outputs
25//! - [`Op`]: the operation kind with parameters
26
27pub mod ad;
28pub mod apple_bnns;
29pub mod async_copy;
30pub mod attention_layout;
31pub mod audio;
32pub mod capability;
33pub mod const_check;
34pub mod dtype;
35pub mod dynamic;
36pub mod env;
37pub mod env_catalog;
38pub mod env_registry;
39pub mod equivalence;
40pub mod fft;
41pub mod graph;
42pub mod hir;
43pub mod infer;
44pub mod infer_shape;
45pub mod inspect;
46pub mod layout;
47pub mod lir;
48pub mod logical_kernel;
49pub mod lowp_codec;
50pub mod measure;
51pub mod mir;
52pub mod module;
53pub mod numeric_check;
54pub mod nvfp4;
55pub mod op;
56pub mod op_registry;
57pub mod opcodes;
58pub mod ops;
59pub mod perfetto;
60pub mod phase;
61pub mod pretty;
62pub mod provenance;
63pub mod quant;
64pub mod region_encode;
65pub mod residual;
66pub mod rng;
67pub mod tensor_inspect;
68pub use nvfp4::{FP4_E2M1_LUT, NVFP4_GROUP_SIZE, fp4_e2m1_to_f32, fp8_e4m3_scale_to_f32};
69pub mod binding_manifest;
70pub mod component;
71pub mod hir_extension;
72pub mod reflect;
73pub mod rf;
74#[cfg(feature = "serialize")]
75pub mod serialize;
76pub mod shape;
77pub mod target;
78pub mod text;
79pub mod variant;
80pub mod verify;
81
82pub use ad::AdPipelineStage;
83pub use async_copy::{AsyncCopy, BarrierToken, DoubleBuffer, SyncCopy};
84pub use attention_layout::{
85    ATTENTION_FLASH_MAX_HEAD_DIM, AttentionGeom, AttentionLaunchStrides,
86    attention_dispatch_use_row, attention_geom, attention_launch_strides, cpu_attention_bshd,
87    cpu_attention_packed_bshd_qkv, detect_packed_bshd_qkv_attention, mask_strides_bhsd,
88    mask_strides_for_shape, packed_bshd_narrow_elidable, packed_bshd_qkv_strides, strides_bhsd,
89    strides_bshd, strides_for_shape,
90};
91pub use capability::{N_KINDS, OpCaps, fused_kinds};
92pub use dtype::{DType, Element, ElementSubtype, scalar_constant_bytes};
93pub use dynamic::sym;
94pub use dynamic::{
95    DimEnv, bind_graph, collect_dynamic_symbols, has_dynamic_dims, infer_bindings_from_f32_inputs,
96    infer_bindings_from_inputs, same_binding, sync_concat_shapes, sync_expand_ops,
97    sync_graph_shapes, sync_narrow_ops, sync_reshape_ops,
98};
99pub use env::{RlxEnv, RuntimeOverrides, flag, is_unset, parse_or, set, unset, var, var_os};
100pub use env_catalog::{EnvVarDoc, format_catalog as format_env_catalog};
101pub use env_registry::{
102    EnvKind, EnvLayer, EnvStability, EnvVarEntry, REGISTRY as ENV_REGISTRY, canonical_name,
103    format_registry_markdown, is_registered, lookup as env_lookup, public_catalog_docs,
104};
105
106/// Public `RLX_*` catalog (`just env-catalog`).
107pub fn env_catalog() -> &'static [EnvVarDoc] {
108    public_catalog_docs()
109}
110
111pub use equivalence::{
112    IgnoreConfig, fingerprint, node_value_key, ops_deep_eq, structural_hash, structurally_eq,
113};
114pub use fft::{FftGpuPlan, FftMeta, FftNorm, fft_meta, fftn_axes_all, normalize_fftn_axes};
115pub use graph::{Graph, Node, NodeId};
116pub use hir::{
117    FusionPolicy, HirGraphExt, HirModule, HirMut, HirNode, HirNodeId, HirOp,
118    window_token_gather_bsn, window_token_scatter_bsn,
119};
120pub use infer::GraphExt;
121pub use inspect::{
122    inspect_buffer_plan, inspect_graph, inspect_graph_diff, inspect_hir, inspect_hir_stats,
123    inspect_lir, inspect_mir, inspect_mir_diff, inspect_mir_stats,
124};
125pub use layout::{Coord2, Ragged, ShapeTuple, Strides2, Strides3, Tile2, Tile3};
126pub use lir::{
127    LirBufferPlan, LirBufferSlot, LirFingerprint, LirIoManifest, LirModule, LirViewAlias,
128};
129pub use logical_kernel::{
130    KernelDispatchConfig, KernelDispatchPolicy, LogicalKernelEntry, logical_kinds_in_graph,
131    registered_logical_kernels, should_lower_to_common,
132};
133pub use measure::{CacheBuster, Tick, time_ns};
134pub use mir::{MirModule, MirNode, MirNodeId, MirOp};
135pub use module::{GraphModule, GraphStage};
136pub use numeric_check::{BadValue, DebugMode, DebugScanner, NanReport, check_node, first_bad};
137pub use op::{
138    AdaNormKind, ChainOperand, ChainStep, Op, OpKind, PadMode, RegionPrologue, RopeStyle,
139    ScatterNdReduction, SpdMatFn, SynthBwdWrt, SynthKind, TransformStep,
140};
141pub use op_registry::{
142    CustomOpError, JvpContext, LowerContext, OpExtension, OpRegistry, VjpContext, VmapContext,
143    global_registry, is_op_registered, lookup_op, register_op, register_op_strict,
144};
145pub use ops::attention::attention_kind_op;
146pub use ops::backward::{unpack_ada_layer_norm_backward, unpack_gated_residual_backward};
147pub use ops::dsp::{FirMode, iir_impulse_response};
148pub use phase::{Phase, PhaseSchedule, derive_phases};
149pub use provenance::{NodeOrigin, node_label, stamp_pass_origins};
150/// Alias kept for existing `ENV_CATALOG` imports.
151pub use public_catalog_docs as ENV_CATALOG;
152pub use quant::{QuantMap, QuantScheme, ScaleLayout, ScaledFormat};
153pub use region_encode::{
154    FK_BATCH_SINGLE_KERNEL_MAX, PrologueLaunchGrid, REGION_META_WORDS, REGION_PROLOGUE_NONE,
155    REGION_PROLOGUE_RESIZE_NEAREST_2X_NCHW, RegionNchwDims, batch_region_slice_dst_off_f32,
156    batch_region_slice_elems, batch_region_slice_shape, encode_chain_operand, encode_chain_steps,
157    encode_elementwise_region_meta, encode_prologue_tail, fk_batch_single_kernel_enabled,
158    fk_batch_use_single_launch,
159};
160pub use rng::{
161    BNNS_AES_CTR_STATE_SIZE, BnnsAesCtr128, GUMBEL_UNIFORM_EPS, NativeBnnsGumbelMax, Philox4x32,
162    RngBackend, RngOptions, argmax_f32, bnns_categorical_from_uniform, bnns_entropy_seed,
163    bnns_generator_state_is_fresh, bnns_gumbel_from_uniform, bnns_gumbel01_from_bytes,
164    bnns_seed_from_generator_state, bnns_uniform_f32_from_bytes, bnns_uniform_f32_from_u32,
165    bnns_uniform_int_from_u32, combine_seed, fill_normal_like, fill_uniform_like,
166    gumbel_max_argmax_with, gumbel01_from_uniform, ort_engine_seed,
167};
168#[cfg(feature = "serialize")]
169pub use serialize::{
170    LIR_BIN_VERSION, hir_from_json, hir_to_json, lir_from_bytes, lir_from_json, lir_to_bytes,
171    lir_to_json,
172};
173pub use verify::{VerifyError, verify, verify_all, verify_shapes};
174
175/// Lower a HIR module to MIR, then extract the legacy [`Graph`] API surface.
176pub fn hir_to_graph(hir: HirModule) -> Result<Graph, hir::LowerError> {
177    Ok(hir.lower_to_mir()?.into_graph())
178}
179pub use binding_manifest::{BindingManifest, IoBindingEntry, WeightBlock};
180pub use component::{CompilationMode, ModelComponent};
181pub use hir_extension::{
182    HirExtensionFn, apply_hir_extensions, apply_hir_extensions_named, register_hir_extension,
183    registered_hir_extensions,
184};
185pub use reflect::{
186    BlockSpecialization, HirReflection, ManifestDiff, MirReflection, SpecializeBlockRecord,
187    layout_for_binding, layout_from_lir, probe_block_specialization, symbolic_layout_hint,
188};
189pub use rf::{
190    complex_div, const_f32, cs_degen_z_in, find_param_node, find_param_nodes, mag2, s11_from_z,
191    scalar_f32,
192};
193pub use shape::{Dim, DimBinding, Shape, ada_modulation_launch, ada_modulation_lead_pack};
194pub use variant::{ModelPhase, ModelVariant};