[RFC]: Native Laya inference with Rust and CUDA

[linear3735]

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.so bundle. 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.

[xiaoyu-xyz]

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.py and
tools/build.py together, the current split is narrower than it looks:

• export.py runs TileLang and writes a standalone generated.cu
  (:262) plus a manifest.json carrying tilelang_version and
  nvcc_flags (:263-274).
• build.py compiles only .cu files — generated.cu, kernels/runtime.cu,
  kernels/model_ops.cu — with nvcc (:23-27).

So TileLang is needed to generate the kernel, and its output is ordinary
CUDA that nvcc compiles on its own. The generated source is the artifact; the
generator is not.

Proposal

Keep TileLang as the generator, and take it out of the build path:

1. Check in generated.cu alongside the hand-written runtime.cu and
   model_ops.cu, together with the provenance manifest (TileLang version,
   nvcc flags, source hashes). It is ~generated code, reviewed like any other
   change, and export.py remains the way to regenerate it.
2. build.sh compiles only .cu files, the way cua_s1: add a native CUDA text worker #19's does, so building the
   bundle needs nvcc and nothing else.
3. Declare the artifact in a manifest so CI can find it, per the contract
   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

1. Reproducibility. A binary built from a checked-in .cu is reproducible
   from nvcc plus the manifest. The same binary built through TileLang depends
   on the TileLang version, its CUTLASS include tree, and its TVM host-code
   generation, all of which move independently of this repository. The
   build-manifest.json you already emit is the right instinct; checking in the
   source makes it verifiable rather than only recorded.
2. CI can compile it. The compile job needs a CUDA toolkit image and no
   Python environment. As it stands, compiling this bundle in CI requires
   installing tilelang and its dependencies, which is why the validation
   report'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-features to the Rust steps; no workflow step compiles CUDA.
3. One compile path for every backend. cua_s1: add a native CUDA text worker #19's qwen3_5/build.sh already has
   the 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.py
   gets 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.json declaring sources, the build script, the output library
  name, the architectures actually built, and (once measured) the tolerance.
• ./build.sh <output-dir> <compute-capability> producing exactly the declared
  library, executable, with the architecture list visible in the script.
• A validated backend also naming a reference entrypoint that a GPU runner can
  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,
int returns, work queued on a caller-supplied stream for Graph capture, tuned
GEMM 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 are
built for sm_89 with sm_80-level instructions. sm_90a is architecture-locked
and 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_90a as the optimised target but make sure a portable
baseline (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#1 adds build.sh and laya.backend.json to this
  branch — a wrapper over export.py and build.py that reaches the interface
  below, 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.
• cuda: add the backend contract and its checker #25 adds the contract document and the checker that
  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 in
generated.cu and compiling only that with nvcc is the separate change I
described above, and it is yours to make.