Your PyTorch and JAX, running on Loom
Loom takes Torch/Triton and JAX/Pallas through one compiler and one IR pipeline. Use it as a live backend that executes your framework and hands back a native array, or ahead of time, compiling a model into an object the chip runs on its own. The user describes intent. The object describes exactly what the chip will do.
Same compiler, same IR, three entry points
Run stock Torch/Triton or JAX/Pallas with Loom as the execution device. The compiler lowers your code and the chip executes it in the request path, returning a native torch.Tensor or jax.Array. This is the familiar accelerator pattern: Loom in the slot a GPU would occupy, without a host round-trip in the inner loop. It works today and the op coverage is still growing.
Point the compiler at the same framework code and, instead of executing live, get a .loom object file: a deterministic, inspectable artifact you version, diff, replay, and run with bounded timing. The object file is the contract. This is the path for deployments where predictable execution matters more than dropping into an existing Python process.
Skip framework code entirely. loomc ingest-checkpoint takes a pretrained model, a Hugging Face checkpoint directory with config.json and .safetensors (sharded or single file), and lowers it to a .loom object file that the on-chip ACP runs on its own. Use original floating-point weights, apply one Loom quantization step, and get a manifest that states exactly what happened. Bring a model, get an object file that runs host-free.
One internal truth, regardless of frontend
Loom has four frontends and one internal compiler. Every frontend funnels through the same lowering stack: the same Tile IR, the same CFG IR, and the same packaged .loom object file. The .loom object file contains sections such as Tile IR, CFG IR, command rows, runtime dispatch payloads, and validation and debug reports. Compiler decisions live in one place.
Loom accepts four ingress paths: ONNX, Torch FX, JAX/Pallas, and Triton as a constrained tile-oriented frontend. All four lower into the same Loom Tile IR and CFG IR. JAX/Pallas exposes the Loom contract most directly; Triton exposes a supported subset shaped by Loom hardware semantics; ONNX and FX reach the contract through import and lowering.
Frontend code lowers into a frontend-neutral Tile IR representing tile-local compute and movement. The IR carries Loom-native primitives directly: mxu.gemm, sce.scan, sce.fft, dma.prefetch, rom.lookup, acp.launch, and the effect/token surface that makes pipelines and events explicit. No pretending everything is generic elementwise math.
The Tile IR lowers into a control flow graph IR that owns control topology and the physical realities of the hardware: blocks, loops, and joins; buffer placement across SPAD0 and SPAD1; memory lifetimes and interference; double buffering; pipeline stages; and dependency edges between MXU, SCE, DMA, and ACP. This is where the compiler commits to what the chip will actually do.
The compiler resolves padding, layout fixups, and overlap planning. Shallow issue-group scheduling produces deterministic engine overlap without exposing raw VLIW to users. Legality checks surface problems early, and fallback reports say exactly which intrinsic was chosen and why.
Graph import, kernel frontends, and a native DSL, all into one Tile IR
ONNX is a portable interchange from any training framework. For PyTorch, register the backend once and torch.compile(model, backend="loom") runs the model on Loom, handing back a native torch.Tensor. A separate path lowers Torch modules to .loom objects without live execution.
An XLA/PJRT plugin runs stock JAX on Loom with no Loom-specific rewrite: an ordinary jax.jit returns a native jax.Array, and a standard pallas_call grid maps natively to the hardware.
Loom ingests structured upstream Triton and TorchInductor TTIR/TTGIR directly; unsupported GPU-specific forms fail closed with precise diagnostics. loomc builds .loom objects from TTIR or a TorchInductor cache. For native authoring in the same style, LTL (loom.triton.language) preserves the Triton mental model while exposing Loom hardware controls directly.
A Pallas-like DSL for hand-tuned control of DMA, SRAM windows, and block grids, when you want to decide the data movement a general-purpose compiler would decide for you.
Full detail, code, and tooling: Frontends.
loomc: The CompilerTile IR to CFG IR to packaged object file
loomc owns placement, DMA overlap, engine handoff, event and barrier synthesis, scheduling, bundling, and object-file packaging. The user describes tile math and intended dataflow. The compiler turns that into fused, pre-scheduled command streams that keep every engine fed. This is what makes Loom tractable to use without exposing raw cmd64 choreography, SPAD packing, or event-bit reservation.
When a target SKU exposes ROM-backed intrinsics, the compiler selects them automatically. The same source lowers to a programmable command stream on standard Loom and to a Structured Overlay Package on Loom-ROM, without forking the toolchain or forcing a second platform.
Data movement is the real cost on any accelerator. By default, loomc plans SRAM placement across MXU and SCE engines, double-buffers where it helps, and fuses adjacent operations into single command sequences so that intermediate tensors never leave the chip unnecessarily. A kernel author who wants direct control drops to the ltl.* dialect and stages SRAM and DMA explicitly; the same legality and fusion guarantees still apply.
Every significant compiler decision is inspectable. loomc explain surfaces tiling choices, placement rationale, fallback reports, and which Loom-native intrinsics were actually lowered. Autotune picks are recorded with their trade-offs. The compiler does not hide its reasoning behind a black box; the validation and debug sections of the .loom object file spell it out.
libloom: The RuntimeKeeping the hot path fed, staged, and fused
libloom handles device discovery, capability readout, buffer management, command submission, SRAM allocation, double-buffering, and weight-prefetch scheduling, keeping the hot path fed, staged, and fused.
Develop before the silicon arrives. libloom exposes a unified surface that spans file-backed hardware emulation and the production /dev/loom kernel path. Application teams build pipelines, exercise the runtime, and prove bit-exact backend parity in emulation, then drop the same code onto a physical Loom Node with zero modifications. The runtime abstracts the PCIe submission queues; the API contract never changes.
The runtime is a C11 core with a versioned, stable ABI. C++, Rust, and Python bindings are thin layers over that same ABI. A model deployed in C behaves identically when driven from a C++ control application, a Rust service binary, or a Python bring-up notebook.
For large models that do not fit in SRAM, libloom stages weights from external memory in the background while compute engines process the current layer. The runtime coordinates with the compiler's placement plan so prefetch never fights the critical path.
Common workload patterns are exposed as named, stable helper surfaces. Attention and paged-KV gather/scatter, vision patch extraction and score-topk, norm/activation epilogues, masked-softmax sequences, overlap-save spectral pipelines, LoRA adapter fusion and AdamW microsteps, replayable random fills, sparse shortlist compaction, TurboQuant encode/decode/score, and fixed-rank tensor DMA / snapshot / trace control planes all ship as first-class libloom families. Each one is paired with a runnable reference example and a deterministic test corpus.
Host code submits command buffers through standard queues and receives completion events through writeback, MSI-X, and mailboxes. The runtime abstracts the submission model, and the same API works across every Loom SKU and derivative.
loomctl: Bring-up and DiagnosticsInspectable, replayable, debuggable execution
loomctl provides device introspection, ROM-table discovery, device health monitoring, and low-level diagnostics. It is the first tool developers reach for when a board comes up, and the one they keep running when production silicon needs to be observed in the field.
Versioned buffers, event signaling, completion writeback, and trace markers make execution inspectable and replayable. Teams can debug pipelines, verify fallbacks, and see what actually ran.
Combined with Loom's reproducible sampling support, loomctl can capture a snapshot of working state and replay branches deterministically. That turns debugging from a forensic exercise into a reproducible one, which matters for safety cases, regression analysis, and release gating.
Submit once, let the device run the loop
For workloads like decode, state-space inference, and other long-lived local loops, the host can submit work once and let the ACP handle token loops, stop conditions, retries, threshold checks, and multi-engine dispatch on the device. That removes per-step host round-trips from the critical path.
The ACP runs loops with cycle budgets, fault conditions, and fallback paths defined at submission time. If a loop misses its budget or hits a fault condition, the device falls back deterministically and raises an event.
The host and ACP communicate through mailboxes, event bits, and trace markers. Host software can poll, wait on interrupts, or subscribe to specific events, depending on the application's latency posture.
Curated, deterministic patterns across inference, signal, control, and board scale
The Loom reference workload library is organized around deterministic patterns paired with test corpora:
Direct backend ops and the Loom-ROM path
When a kernel author needs to step below the standard intrinsic surface, loomc exposes an inline_backend_op override. Direct backend ops sit inside the same Tile IR and .loom object file as normal kernels; they are explicitly marked non-default in the object-file sections so maintainers can see what opted out of compiler-owned scheduling and why. This override is for the narrow case where a specific fusion or a non-standard quantization flow is required before the compiler roadmap catches up. It is not the default path.
The same host software model extends into Loom-ROM. In the fixed-model path, loomc emits a Structured Overlay Package, and standard Loom software feature-detects the ROM-backed variant at runtime. One SDK, one runtime, one set of tools, from programmable Loom through fixed-function derivatives.