Keep Laya weights and workspace on CUDA

Cargo.toml

@@ -1,3 +1,3 @@
 [workspace]
-members = ["src/frontend", "src/models/cua_s1/native", "src/models/laya"]
+members = ["src/frontend", "src/models/cua_s1/native", "src/models/laya", "src/backends/cuda"]
 resolver = "3"

recipe/laya/README.md

@@ -56,3 +56,52 @@ if the worker requires a bearer token.
 
 See the [frontend documentation](../../src/frontend/README.md) for configuration
 and transport behavior.
+
+## Native residency and workspace validation
+
+These opt-in Rust checks test CUDA allocations and transfers. They do not need the
+Python worker or frontend and do not test inference, model outputs or latency.
+The normal CPU tests skip them.
+
+Use a Linux host with an approved CUDA GPU, a working NVIDIA driver, the CUDA
+toolkit (`nvcc`) and Rust. Build the trusted resource library from this checkout;
+it uses ABI version 1 and needs neither TileLang nor cuBLAS. `LAYA_CUDA_DEVICE` is
+the approved device ordinal after `CUDA_VISIBLE_DEVICES` filtering.
+
+```sh
+export LAYA_CUDA_LIBRARY=/tmp/liblaya-resources.so
+export LAYA_CUDA_DEVICE=0
+nvcc -shared -Xcompiler=-fPIC -O2 src/backends/cuda/kernels/runtime.cu \
+  -o "$LAYA_CUDA_LIBRARY"
+```
+
+The workspace check needs no checkpoint. It writes and reads all 17 buffers twice
+at `(batch, sequence) = (1, 16), (1, 512), (16, 512)`, using deterministic byte
+patterns:
+
+```sh
+cargo test --release --locked -p omni-laya --lib \
+  workspace::tests::real_gpu_workspace_capacity_and_reuse \
+  -- --ignored --exact --nocapture
+```
+
+For the residency check, use an unchanged local snapshot of
+`convaiinnovations/laya` revision `55cf4c4ebb4ebe31b2550e8bdf3bd21b99753851`,
+including `model.safetensors`. Generate the oracle from that same snapshot in a
+Python environment with PyTorch, safetensors and NumPy:
+
+```sh
+export LAYA_CHECKPOINT=/path/to/laya/snapshot
+export LAYA_WEIGHT_ORACLE=/tmp/laya-weight-oracle.json
+python recipe/laya/native/export_weights.py "$LAYA_CHECKPOINT" "$LAYA_WEIGHT_ORACLE"
+cargo test --release --locked -p omni-laya --lib \
+  resident::tests::real_checkpoint_residency_matches_torch \
+  -- --ignored --exact --nocapture
+```
+
+The residency check validates all 206 checkpoint tensors, uploads the 205 used
+tensors and compares readback hashes with the Torch conversion oracle. The legacy
+`temperature` buffer is validated but not uploaded. Reported allocation bytes
+exclude CUDA context and library overhead. See the
+[model contracts](https://github.com/linear3735/system1-omni/blob/codex/laya-workspace/src/models/laya/README.md) for storage precision,
+workspace layouts and ownership.

src/backends/cuda/Cargo.toml

@@ -0,0 +1,16 @@
+[package]
+name = "omni-cuda"
+version = "0.1.0"
+edition = "2024"
+publish = false
+
+[dependencies]
+anyhow = "1"
+libloading = "0.8"
+
+[dev-dependencies]
+tempfile = "3"
+
+[[test]]
+name = "runtime"
+path = "../../../tests/backends/cuda/runtime.rs"

src/backends/cuda/README.md

@@ -1,7 +1,56 @@
 # CUDA backend
 
-Planned home for high-performance NVIDIA GPU operations and kernel integration. Implement the operations required by the first model, with hardware-specific optimizations where needed.
+[`qwen3_5/`](qwen3_5/) provides the prefill-only Qwen3.5 operations used by the
+Cua-S1 native worker, measured on sm_89.
 
-Model orchestration, batching policy, state management, and kernel selection remain with the model engine. CUDA and Metal implementations do not need identical internal structures or a universal tensor abstraction.
+## Laya resources
 
-Status: [`qwen3_5/`](qwen3_5/) has the operations of a prefill-only Qwen3.5 forward pass, used by the Cua-S1 native worker and measured on sm_89. Other models are planned.
+`omni-cuda` loads Laya's CUDA resource library at runtime. It owns one device and
+stream per context, plus the buffers allocated through that context. Rust builds
+and CPU tests need no CUDA toolkit.
+
+This first slice covers allocation, copies, synchronization and cleanup. Model
+initialization, weights, kernel calls, Graphs and hardware-specific optimizations
+remain separate. It does not yet run Laya inference.
+
+### Build and check
+
+On a machine with the CUDA toolkit, build the resource library:
+
+```sh
+nvcc -shared -Xcompiler=-fPIC -O2 src/backends/cuda/kernels/runtime.cu -o /tmp/liblaya_cuda.so
+LAYA_CUDA_LIBRARY=/tmp/liblaya_cuda.so LAYA_CUDA_DEVICE=0 \
+  cargo test --locked -p omni-cuda --test runtime -- --ignored
+```
+
+The device is an ordinal after `CUDA_VISIBLE_DEVICES` filtering. The library
+contains no generated kernels and needs neither TileLang nor cuBLAS. This command
+builds only the resource slice; the complete model bundle has a separate build.
+
+The normal CPU tests compile a small C fixture with `cc`. They check the dynamic
+loader, errors, copy bounds and resource lifetime. They do not validate CUDA or
+hardware support. The ignored test exercises real allocation and copy roundtrips.
+
+### Ownership and ABI
+
+Load only a trusted library with the matching ABI. `Cuda::load(path, device)`
+checks `laya_abi_version() == 1` and all required symbols before creating a stream.
+The old prototype's `laya_init` library has no version symbol and is rejected.
+
+`Cuda` and `Buffer` stay on their creating thread. A buffer keeps its stream and
+library alive even after the caller drops `Cuda`. Operations select the owning
+device before using its resources. Destruction attempts synchronization and
+cleanup; call `sync()` explicitly when errors need to reach the caller.
+
+`write` and `read` check byte limits and synchronize before returning, so borrowed
+host memory cannot outlive a queued copy. They are not Graph-capture operations.
+Allocation of zero bytes is rejected; empty reads and writes are no-ops.
+
+The native resource entry points return zero on success and CUDA error codes on
+failure; code 1000 means an invalid runtime argument. `laya_error_string` explains
+the code. Upload and download take the caller's stream as their last argument and
+do not synchronize internally. No Hopper requirement or model initialization is
+hidden in stream creation.
+
+These are Laya's resource entry points, not a new shared tensor interface. A common
+runtime can be extracted when another model needs the same implementation.

src/backends/cuda/kernels/runtime.cu

@@ -0,0 +1,65 @@
+#include <cuda_runtime.h>
+#include <stdint.h>
+
+namespace {
+constexpr int invalid_argument = 1000;
+}
+
+extern "C" {
+uint32_t laya_abi_version() { return 1; }
+
+const char* laya_error_string(int code) {
+  return code == invalid_argument ? "invalid runtime argument"
+                                 : cudaGetErrorString(static_cast<cudaError_t>(code));
+}
+
+int laya_set_device(int device) { return cudaSetDevice(device); }
+
+int laya_stream_create(void** stream) {
+  if (!stream) return invalid_argument;
+  *stream = nullptr;
+  cudaStream_t created = nullptr;
+  cudaError_t status = cudaStreamCreateWithFlags(&created, cudaStreamNonBlocking);
+  if (status == cudaSuccess) *stream = created;
+  return status;
+}
+
+int laya_alloc(void** p, size_t bytes) {
+  if (!p) return invalid_argument;
+  *p = nullptr;
+  if (!bytes) return invalid_argument;
+  void* allocated = nullptr;
+  cudaError_t status = cudaMalloc(&allocated, bytes);
+  if (status == cudaSuccess) *p = allocated;
+  return status;
+}
+
+int laya_free(void* p) {
+  if (!p) return invalid_argument;
+  return cudaFree(p);
+}
+
+int laya_upload(void* dst, const void* src, size_t bytes, void* stream) {
+  if (!bytes) return 0;
+  if (!dst || !src || !stream) return invalid_argument;
+  return cudaMemcpyAsync(dst, src, bytes, cudaMemcpyHostToDevice,
+                         static_cast<cudaStream_t>(stream));
+}
+
+int laya_download(void* dst, const void* src, size_t bytes, void* stream) {
+  if (!bytes) return 0;
+  if (!dst || !src || !stream) return invalid_argument;
+  return cudaMemcpyAsync(dst, src, bytes, cudaMemcpyDeviceToHost,
+                         static_cast<cudaStream_t>(stream));
+}
+
+int laya_sync(void* stream) {
+  if (!stream) return invalid_argument;
+  return cudaStreamSynchronize(static_cast<cudaStream_t>(stream));
+}
+
+int laya_stream_free(void* stream) {
+  if (!stream) return invalid_argument;
+  return cudaStreamDestroy(static_cast<cudaStream_t>(stream));
+}
+}