Repository navigation
[RFC]: Native Laya inference with Rust and CUDA #14
Description
Motivation
#2 added the Rust frontend with an external model worker. This proposal adds an in-repository Laya engine that loads the original weights and runs inference directly through CUDA.
Delivery
Implementation will land in small PRs, with at most 500 core changed lines per PR. #21 starts with CPU checkpoint loading and validation. Input processing, model execution, CUDA optimizations and HTTP integration follow as separate changes. The full prototype from #16 remains available for reference.
Design
Start with Laya 0.3.20's English model and the BF16 fast path on Hopper. Support choice, score, and noul through the existing HTTP API.
HTTP → Rust preprocessing → GPU worker → CUDA inference → Rust response
Follow the layout in #6:
| Directory | Responsibility |
|---|---|
src/frontend/ |
HTTP, request lifecycle, and the small engine interface |
src/models/laya/ |
Tokenizer, weights, model execution, and output decoding |
src/backends/cuda/ |
CUDA resources, Graphs, and kernels |
recipe/laya/native/ |
Build, run, and validation scripts |
- One GPU worker processes requests sequentially. The queue holds 32 requests and returns HTTP 503 when full. A 30-second deadline includes upload and queueing. Readiness follows loading and warmup. There is no cross-request batching.
- CUDA Graphs cover the encoder and decision transformer. A new shape allocates buffers, runs two warmups, and captures synchronously. An LRU cache retains up to four shapes and 512 MiB of workspaces. The variable-size scorer stays outside Graph.
- Use the optimized RoPE kernel, with separate QKV, RoPE, and attention launches.
- Python and TileLang generate CUDA sources; nvcc builds a reusable
liblaya_cuda.sobundle. Compiled binaries are not checked in. Deployment uses the Rust binary, checkpoint, bundle, and CUDA/cuBLAS libraries.
Implementation and checks
Implementation · Build and validation · Results and samples
Checks cover tokenizer and weight parity, eager/Graph outputs, and HTTP concurrency and lifecycle. Boundary cases produced the same decisions with a maximum numeric difference of 0.0014 against a 0.002 limit.
H800 BF16 measurements on 2026-09-27: warmed short-request engine wall time was 2.84 ms for official fast and 1.81 ms for native, both using the same RoPE. Native HTTP C1 median was 2.00 ms. The validation report records the measured source hashes and timing method.
Activity
cuda: add the backend contract, a contract check, and CUDA CI
cuda: add the compile-matrix and parity drivers
xiaoyu-xyz commentedon Sep 28, 2026
I would like to work on the build side of this, and I have a concrete
suggestion that I think removes a disagreement rather than picking a side in one.
The question in #3 was never answered
In #3 I asked whether the backend should be CUDA C++ or TileLang. It was not
settled, and both answers are now in the tree: this issue generates CUDA from
TileLang, while #19 hand-writes CUDA C++ with cuBLASLt and states "no mistral.rs,
candle or Triton underneath". #12 plans Triton. That is three build systems
before a second model has landed, and #9 asks for CUDA paths for further models.
I do not think the useful question is "which one". Reading
tools/export.pyandtools/build.pytogether, the current split is narrower than it looks:export.pyruns TileLang and writes a standalonegenerated.cu(
:262) plus amanifest.jsoncarryingtilelang_versionandnvcc_flags(:263-274).build.pycompiles only.cufiles —generated.cu,kernels/runtime.cu,kernels/model_ops.cu— withnvcc(:23-27).So TileLang is needed to generate the kernel, and its output is ordinary
CUDA that
nvcccompiles on its own. The generated source is the artifact; thegenerator is not.
Proposal
Keep TileLang as the generator, and take it out of the build path:
generated.cualongside the hand-writtenruntime.cuandmodel_ops.cu, together with the provenance manifest (TileLang version,nvcc flags, source hashes). It is ~generated code, reviewed like any other
change, and
export.pyremains the way to regenerate it.build.shcompiles only.cufiles, the way cua_s1: add a native CUDA text worker #19's does, so building thebundle needs
nvccand nothing else.below.
This keeps every argument for TileLang that I can see — the fused GEMM epilogues
and the selected RoPE geometry stay exactly as they are, and nobody rewrites a
working kernel in another language — while removing TileLang, Python and TVM
from the set of things needed to reproduce a binary.
Why the build path matters more than the kernel path
.cuis reproduciblefrom
nvccplus the manifest. The same binary built through TileLang dependson the TileLang version, its CUTLASS include tree, and its TVM host-code
generation, all of which move independently of this repository. The
build-manifest.jsonyou already emit is the right instinct; checking in thesource makes it verifiable rather than only recorded.
Python environment. As it stands, compiling this bundle in CI requires
installing
tilelangand its dependencies, which is why the validationreport's last line — "The GitHub CPU workflow checks all Rust features. It
does not compile CUDA kernels or validate model execution, numerical parity,
or GPU latency" — is still true after Add native Laya inference with Rust and CUDA #16. That PR only added
--all-featuresto the Rust steps; no workflow step compiles CUDA.qwen3_5/build.shalready hasthe shape CI wants: an output directory and a compute capability. If this
backend exposes the same two-argument script, one CI job compiles both, and
a third model reuses it unchanged. A wrapper over
export.py+build.pygets there without changing either program.
The interface I would check both backends against
I opened no PR for this yet because it is your call as the author of this
design. Concretely, I would have CI verify, per backend directory:
<name>.backend.jsondeclaring sources, the build script, the output libraryname, the architectures actually built, and (once measured) the tolerance.
./build.sh <output-dir> <compute-capability>producing exactly the declaredlibrary, executable, with the architecture list visible in the script.
execute for parity.
That is deliberately small: it says nothing about how kernels are written. The
matching question in #19's terms is whether
ops.h's C ABI —*_abi_version,intreturns, work queued on a caller-supplied stream for Graph capture, tunedGEMM plans imported rather than re-tuned — should be the common surface. I think
it should, because it is the part that lets one model's engine call another
model's kernels, which is what makes sharing possible at all.
One detail worth deciding now
This bundle is pinned to
compute_90a(export.py:268) and #19's kernels arebuilt for
sm_89with sm_80-level instructions.sm_90ais architecture-lockedand will not run on Ada, and it is not forward-compatible with later
architectures. If the goal is one backend directory that several models can use,
I would keep
sm_90aas the optimised target but make sure a portablebaseline (PTX, or sm_80-and-later machine code as #19 does) is what CI compiles
and what a new contributor without a Hopper card can run. Reproducing your H800
numbers should not be a prerequisite for contributing a kernel.
Two things are already open, so this is a proposal with code behind it rather
than a plan:
linear3735/system1-omni#1addsbuild.shandlaya.backend.jsonto thisbranch — a wrapper over
export.pyandbuild.pythat reaches the interfacebelow, without touching either of them. It refuses a compute capability other
than 90 rather than ignoring it, because the export is pinned to
compute_90a.defines that interface, independently of any one model.
Neither of them removes TileLang from the build: the wrapper still calls
export.py, so a machine building this bundle still needs it. Checking ingenerated.cuand compiling only that withnvccis the separate change Idescribed above, and it is yours to make.
feat(frontend): add the in-process engine contract and HTTP service
feat(frontend): add the in-process engine contract and HTTP service
feat(frontend): add the in-process engine contract and HTTP service