GPUs and CUDA

The project trains and tests its model on graphics cards (GPUs), not on the laptop's processor. The programs that run on a GPU are called kernels, and the project writes them in CUDA, NVIDIA's language for GPUs. Every kernel is compiled before it runs. For the active model, an audit reads the compiled result and refuses instructions the project has banned: decimal arithmetic, tensor-core matrix instructions and approximate exponentials. The cards are rented by the hour from a marketplace through one script, under hard limits on how many and how much. The rented GPUs are used for qualification runs and a few short training doses. No trained language model has come out of them yet.

Labels used below: proof means derived from definitions or code, measurement means a recorded run, hypothesis means an untested idea. See docs/formal_vocabulary.md and the glossary.

flowchart LR
  subgraph K["What the active kernels compute"]
    READ["read and write one vector (forward)"]
    PROP["score candidate moves (counterfactual)"]
    PICK["choose and apply one move"]
    EVAL["evaluate: top-1 and top-8 hits"]
    CNT["count tokens: frequency baselines"]
  end
  subgraph B["How they reach the GPU"]
    NVCC["nvcc builds PTX ahead of time"]
    AUD{"PTX admission audit"}
    CUBIN["cubin and resource report"]
    LOAD["driver loads the prebuilt module"]
    STOP["build fails"]
  end
  K --> NVCC --> AUD
  AUD -->|admitted| CUBIN --> LOAD
  AUD -->|rejected| STOP

The kernel map. Source files compile to PTX before anything runs, and the audit sits between compiling and loading. There is no run-time compiler.

The kernels and what they compute

The active model is crates/uor-r4-discrete, and its kernels are integer only. The header of discrete.cu says model values use gather, sign, add and subtract, comparison and finite action-table look-ups, and that storage indexing uses integer products. It has 577 lines and 12 entry points:

Five more modules are generated from discrete.cu at build time, each its own source, PTX and cubin, so the earlier images stay byte for byte unchanged: a diagnostic decomposition, baseline reuse, row-parallel scoring (rows, 3 entry points), scoring several coordinates per launch (rowsk, 5) and selecting and committing a set of moves (bvgk, 8). The build also writes exact geometry tables from the core crate: AXES[512], COMPOSE[4096], INVERSE[64], 16 carrier roots, a breadth-first reachability map and 15 generators (build.rs).

Other CUDA code in the repository is not on the active path:

Line counts are raw counts of the files named.

How a kernel is admitted

Compile ahead of time (definition). The build script calls nvcc -ptx -std=c++17 -arch=compute_NN, where NN is the card's capability from the explicit UOR_NVCC_ARCH setting. NVIDIA's run-time compiler, NVRTC, is banned project-wide (D29, amended 13 October): what nvcc compiles at build time is exactly what the driver may load, and banned code present in a module counts even if never launched (D29). Two older pages, the CUDA training runbook and the pods quickstart, still describe NVRTC. They were written before the ban.

The admission audit (proof by inspection of the code). The audit reads the emitted PTX text, not the CUDA source, so a banned instruction cannot hide behind spelling. It strips comments, splits each line into tokens, and rejects the build if any token has a component named f16, f16x2, f32, f64, bf16, bf16x2, tf32, e4m3, e5m2, mma, wmma, wgmma, tcgen05, tensor or tensormap, or starts with ex2.approx. It also requires the set of .entry points to equal the declared inventory exactly: a missing, extra or duplicate entry fails. A self-test runs on every build. It must accept a clean module (floats inside comments are fine) and must reject add.f32, a .bf16 register, mma.sync, ex2.approx, an unexpected entry, a duplicate entry and missing entries. After admission, the same admitted PTX is assembled into a cubin for each of the six modules, and the build fails if the cubin is empty or its resource report is missing.

What it does not say. Integer multiplies are allowed in these kernels, for example to compute indexes. This is a different rule from the serving contract, where no multiply instruction may appear at all (serving and export). The float training crate has no equivalent PTX audit, since float is legal there, and uses precision flags instead (-ftz=false -prec-div=true -prec-sqrt=true, and --fmad=false for the flock selector). Passing the audit shows which instructions are present. It does not show that a kernel computes the right thing. Correctness comes from separate controls on a pod.

Not in continuous integration. No GitHub workflow mentions nvcc or CUDA. The standing checks compile without the cuda feature, and qualification runs on rented pods. For exact-head compile checks there is a recipe, scripts/pod/uor-cuda-check-remote. It rents one fresh RTX 5090 with CUDA 12.8 for 15 minutes (30 for the training crate), runs cargo check and then clippy, and tears the pod down (compute guide). It is compile and lint only, never a model run.

Renting the cards

The rules are written in docs/labs/compute.md and printed by uor-pod itself.

The laptop does non-model work: editing, compiling CPU-only builds, static checks and grading stored outputs. All model compute is GPU work (D24). Serving a finished model on the user's own device is a separate matter and is allowed (charter, section 3.6).

Measured costs.

• On 13 October a verified datacenter RTX 5090 on Vast cost $0.469 an hour, against $1.19 an hour on Runpod. A fresh Vast pod took its toolchain, crate cache and binaries from the private store and was ready in 57 seconds, with nothing built (measurement, compute guide).
• The 19 October A8 dose ran on a Vast pod with two RTX 5090s at $1.816 an hour for about 62 minutes: about $1.9. That is a resource estimate, not an invoice, and storage and transfer were not measured. The owner typed the release for each arm personally (record).
• Scoring 16 candidate moves in one launch took a median 123.57 ms, against 125.51 ms for scoring one. All 5,130 coordinates matched the reference exactly, and a deliberately broken copy failed as it should. The earlier launch used only 8 of the card's 170 multiprocessors (record). Measurement, at that one shape.
• One full pass over every parameter of the discrete model is projected at about 17.8 hours for the route read and 19.2 for the flock read, roughly $16 to $17. That is an unverified projection, not a run (record).
• A fused memory kernel claimed to be 12 times faster was never run on the pod, and its commit broke the build and was reverted. No measurement survives (record).

Where it stands

Go deeper

• Mechanism pages: discrete learning, SpiralCore and FBS, Hamilton gate.
• Code: discrete.cu, build.rs, scripts/pod/uor-pod.
• Guides: pod rules, CUDA training runbook (older, pre-dates the NVRTC ban), pods quickstart.
• Next pages: serving and export, data and tokenizer. Model pages: architecture, training.