gpu: a shared op layer over the backends, views as the currency, and a fourth backend to prove it

.github/workflows/ci.yml

@@ -6,20 +6,6 @@ on:
   pull_request:
 
 jobs:
-  # Static, no build needed: is a GPU op real on the same backends everywhere,
-  # or did one backend land ahead of the others and never get caught up (that
-  # is exactly what happened to concat_part/rope before this job existed —
-  # see tools/check_backend_parity.py's own docstring). Deliberate gaps (the
-  # M9 LLM decode fast path, still CUDA-only) are recorded in
-  # tools/backend_parity_allowlist.txt and don't fail this job; an
-  # unallowlisted asymmetry does.
-  backend-parity:
-    name: backend parity (CUDA / Metal / WebGPU)
-    runs-on: ubuntu-latest
-    steps:
-      - uses: actions/checkout@v4
-      - run: python3 tools/check_backend_parity.py
-
   # All three modes — cpu / gpu / auto — on every OS.
   # What each hosted runner actually offers:
   #   - macos-15 (arm64): Metal works through a paravirtual GPU -> a real GPU test
@@ -40,6 +26,27 @@ jobs:
       - name: Test (cpu / gpu / auto)
         run: ctest --test-dir build -C Release --output-on-failure
 
+  # The host reference backend (gpu_host.h): a "device" that is the CPU, so the
+  # shared GPU layer — gpu_ops.h, the launch policy, residency, the census — and
+  # the conformance tests run on runners with no GPU, where the plain `test` job
+  # above only ever takes the CPU fallback and never enters them. On macOS it
+  # also checks that asking for it by name wins over Metal.
+  host-backend:
+    name: host reference backend / ${{ matrix.os }}
+    runs-on: ${{ matrix.os }}
+    strategy:
+      fail-fast: false
+      matrix:
+        os: [ubuntu-latest, macos-15]
+    steps:
+      - uses: actions/checkout@v4
+      - name: Configure
+        run: cmake -B build -DCMAKE_BUILD_TYPE=Release -DTENSORLIB_HOST_GPU=ON
+      - name: Build
+        run: cmake --build build --config Release --target tensorlib_test
+      - name: Test (cpu / gpu / auto)
+        run: ctest --test-dir build -C Release --output-on-failure -R '^(cpu|gpu|auto)$'
+
   # Build verification for the CUDA toolchain. Hosted runners have no NVIDIA
   # driver, so the kernels are only compiled and linked. Running the tests here
   # doubles as coverage of the dlopen fallback on a driver-less machine, which

CMakeLists.txt

@@ -10,6 +10,7 @@ set(CMAKE_CXX_STANDARD 23)
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
 
 option(TENSORLIB_CUDA "Build the CUDA backend (needs CUDA Toolkit at build time only)" OFF)
+option(TENSORLIB_HOST_GPU "Select the host reference backend (gpu_host.h): runs the GPU layer and its conformance tests with no GPU" OFF)
 
 add_library(tensorlib INTERFACE)
 target_include_directories(tensorlib INTERFACE ${CMAKE_CURRENT_SOURCE_DIR}/include)
@@ -19,6 +20,10 @@ if(APPLE)
     "-framework Accelerate" "-framework Metal" "-framework Foundation")
 endif()
 
+if(TENSORLIB_HOST_GPU)
+  target_compile_definitions(tensorlib INTERFACE TENSORLIB_HOST_GPU)
+endif()
+
 if(TENSORLIB_CUDA)
   # The CUDA backend does NOT link CUDA — the driver is loaded at runtime
   # (dlopen on Unix, LoadLibrary(nvcuda.dll) on Windows; see cuda.h) and kernels

README.md

@@ -161,11 +161,20 @@ that is derived from the pool's measured wake-up on first use
 (`TL_CPU_MIN_WORK` pins it; `misc/census_pool_latency.cpp` shows the
 arithmetic); below that it runs on the calling thread.
 
-New ops sometimes land on one backend (usually CUDA) before the others catch
-up. `tools/check_backend_parity.py` reports, per op, which of CUDA/Metal/
-WebGPU actually implement it rather than falling back to the CPU oracle, and
-fails (in CI too) if an asymmetry isn't recorded in
-`tools/backend_parity_allowlist.txt` as deliberate.
+The GPU ops are written once, in `gpu_ops.h`, over views (`gpu::span`: a device
+handle and a byte offset) and a kernel ABI every backend shares (`gpu_abi.h`). A
+backend is a device core — memory, a kernel table, one `dispatch` — plus kernel
+source, and the ops it runs its own way as members of its `own` struct; an op a
+backend has no kernel for answers false and the evaluator falls back to the CPU.
+`gpu::census(kernel)` counts launches, so a test can tell the two apart. No
+machine here runs CUDA kernels, so `tools/cuda_trace` records what the CUDA
+backend asks of the driver (kernel, grid, every argument) against a stand-in
+`libcuda` in a Linux container, and `tools/cuda_trace/compare.sh <ref>` diffs
+that across a change to the backend's host side.
+`-DTENSORLIB_HOST_GPU=ON` selects a reference backend whose device is the CPU
+(`gpu_host.h`), which runs that whole layer and its conformance tests with no
+GPU. [docs/backends.md](docs/backends.md) covers the layers, the kernel ABI, and
+how to add an op or a backend.
 
 ### Profiling
 

bench/cuda/check/check_attn64.cpp

@@ -15,7 +15,7 @@
 #ifndef TENSORLIB_CUDA
 #define TENSORLIB_CUDA
 #endif
-#include "cuda.h"
+#include "gpu.h"  // cuda.h plus the shared ops (tl::gpu resolves to cuda here)
 #include "kv_cache.h"
 
 #include <algorithm>
@@ -80,21 +80,21 @@ int main() {
       }
       sync_to_host(kn, true);
       sync_to_host(vn, true);
-      if (!cache.append(kn, vn)) { std::printf("  append failed pos %lld\n", (long long)pos); return 1; }
+      if (!cache.append({kn, 0}, {vn, 0})) { std::printf("  append failed pos %lld\n", (long long)pos); return 1; }
 
       if (ci < sizeof(checks) / sizeof(checks[0]) && step == checks[ci]) {
         ci++;
         for (int64_t i = 0; i < HQ * D; i++) hq[i] = rnd();
         sync_to_host(q, true);
-        cache.attn(q, o, HQ, scale);
+        cache.attn({q, 0}, {o, 0}, HQ, scale);
         flush();
         sync_to_host(o, false);
         // Lockstep guard: attn_dpos (split-KV on the capacity-static grid,
         // *d_pos = pos so ctx matches) must be BIT-identical to attn — the
         // device attn_dpos_chunk twins the host attn_split_count/attn_split_chunk
         // heuristic, and the capacity grid's empty splits combine as exact zeros.
         upload_u32(dp, (unsigned)pos);
-        cache.attn_dpos(q, o2, HQ, dp, scale);
+        cache.attn_dpos({q, 0}, {o2, 0}, HQ, {dp, 0}, scale);
         flush();
         sync_to_host(o2, false);
         if (std::memcmp(ho, ho2, (size_t)HQ * D * 4) != 0) {
@@ -150,7 +150,7 @@ int main() {
 
     tl::kv_cache cache;
     if (!cache.init(HKV, MAXC, D)) { std::printf("  cache init failed\n"); return 1; }
-    cache.prefill(qp, ks, vs, op, T, HQ, scale);
+    cache.prefill({qp, 0}, {ks, 0}, {vs, 0}, {op, 0}, T, HQ, scale);
     flush();
     sync_to_host(op, false);
 
@@ -192,8 +192,8 @@ int main() {
     for (int64_t i = 0; i < HKV * D; i++) { hkn[i] = rnd(); hvn[i] = rnd(); }
     for (int64_t i = 0; i < HQ * D; i++) hq1[i] = rnd();
     sync_to_host(kn, true); sync_to_host(vn, true); sync_to_host(q1, true);
-    cache.append(kn, vn);
-    cache.attn(q1, o1, HQ, scale);
+    cache.append({kn, 0}, {vn, 0});
+    cache.attn({q1, 0}, {o1, 0}, HQ, scale);
     flush();
     sync_to_host(o1, false);
 

bench/cuda/check/check_cuda.cpp

@@ -6,15 +6,15 @@
 #ifndef TENSORLIB_CUDA
 #define TENSORLIB_CUDA  // standalone build; the CMake build passes it as a flag
 #endif
-#include "cuda.h"
+#include "gpu.h"  // the shared ops (tl::gpu) resolve to cuda here
 
 #include <cmath>
 #include <cstdio>
 #include <random>
 #include <vector>
 
 using namespace tl::cuda;
-using kop = tl::metal::kop;
+using kop = tl::gpu::kop;
 
 static int failures = 0;
 static void check(bool ok, const char* what) {
@@ -49,7 +49,8 @@ int main() {
       cb[i] = d(g);
       ref[i] = (ca[i] + cb[i]) * 2.0f + 1.0f;
     }
-    bool launched = tl::cuda::binary(kop::add, a, 0, b, 0, o, 0, n, 2.0f, 1.0f);
+    bool launched =
+        tl::gpu::binary(kop::add, {a, 0}, {b, 0}, {o, 0}, n, 2.0f, 1.0f);
     sync_to_host(o, false);  // D2H the device-written output before host read
     bool match = launched;
     for (int64_t i = 0; i < n; i++)
@@ -72,7 +73,7 @@ int main() {
       ca[i] = d(g) * 4.0f;
       ref[i] = 1.0f / (1.0f + std::exp(-ca[i]));
     }
-    bool launched = tl::cuda::unary(kop::sigmoid, a, 0, o, 0, n, 1.0f, 0.0f);
+    bool launched = tl::gpu::unary(kop::sigmoid, {a, 0}, {o, 0}, n, 1.0f, 0.0f);
     sync_to_host(o, false);  // D2H the device-written output before host read
     bool match = launched;
     for (int64_t i = 0; i < n; i++)
@@ -100,7 +101,7 @@ int main() {
         for (int64_t p = 0; p < k; p++) s += ca[i * k + p] * cb[p * n + j];
         ref[i * n + j] = s * 0.5f;
       }
-    bool launched = gemm(a, 0, k, false, b, 0, n, false, o, 0, m, n, k, 0.5f, 0);
+    bool launched = tl::gpu::gemm({a, 0}, k, false, {b, 0}, n, false, {o, 0}, m, n, k, 0.5f, 0);
     sync_to_host(o, false);  // D2H the device-written output before host read
     bool match = launched;
     for (int64_t i = 0; i < m * n; i++)
@@ -131,7 +132,7 @@ int main() {
         ref[i * n + j] = s;
       }
     // trans_b=true, ldb=k (the col stride of the logical k×n = row stride of bt)
-    bool launched = gemm(a, 0, k, false, b, 0, k, true, o, 0, m, n, k, 1.0f, 0);
+    bool launched = tl::gpu::gemm({a, 0}, k, false, {b, 0}, k, true, {o, 0}, m, n, k, 1.0f, 0);
     sync_to_host(o, false);  // D2H the device-written output before host read
     bool match = launched;
     for (int64_t i = 0; i < m * n; i++)
@@ -159,7 +160,7 @@ int main() {
       }
       ref[r] = s * 0.25f + 3.0f;
     }
-    bool launched = tl::cuda::row_op(kop::row_sum, in, 0, o, 0, rows, cols, 0.25f, 3.0f);
+    bool launched = tl::gpu::row_op(kop::row_sum, {in, 0}, {o, 0}, rows, cols, 0.25f, 3.0f);
     sync_to_host(o, false);  // D2H the device-written output before host read
     bool match = launched;
     for (int64_t r = 0; r < rows; r++)
@@ -187,7 +188,7 @@ int main() {
       for (int64_t c = 0; c < cols; c++)
         ref[r * cols + c] = std::exp(ci[r * cols + c] - mx) / sum;
     }
-    bool launched = tl::cuda::row_op(kop::softmax, in, 0, o, 0, rows, cols, 1.0f, 0.0f);
+    bool launched = tl::gpu::row_op(kop::softmax, {in, 0}, {o, 0}, rows, cols, 1.0f, 0.0f);
     sync_to_host(o, false);  // D2H the device-written output before host read
     bool match = launched;
     for (int64_t i = 0; i < rows * cols; i++)

bench/cuda/check/check_llm_decode.cpp

@@ -205,9 +205,9 @@ vec gpu_step(GpuModel& m, int64_t id, int64_t pos) {
     array k = array::rope(h.dot(L.Wk).reshape({HKV, hd}), pos);
     array v = h.dot(L.Wv);  // [1, HKV*hd] == [HKV, hd] bytes
     q.eval(); k.eval(); v.eval();
-    L.cache.append(k.native(), v.native());
+    L.cache.append(k.device_span(), v.device_span());
     array a_out = array::empty({HQ, hd});
-    L.cache.attn(q.native(), a_out.native(), HQ, SCALE);
+    L.cache.attn(q.device_span(), a_out.device_span(), HQ, SCALE);
     array a = a_out.reshape({1, Dm});
     array x1 = x + a.dot(L.Wo);
     array h2 = array::rmsnorm(x1, L.n2, EPS);

bench/cuda/check/trace_sweep.cpp

@@ -0,0 +1,128 @@
+// What the launch trace (tools/cuda_trace) cannot reach through the test suite
+// and the checkers: the kernels no test selects, and the calls no evaluator
+// makes — above all a non-zero OUTPUT offset, which array.h never passes but
+// the contract allows. It checks nothing. It exists to be traced, so that a
+// change to cuda.h's host side shows up in a diff on a machine with no GPU.
+#ifndef TENSORLIB_CUDA
+#define TENSORLIB_CUDA
+#endif
+#include "gpu.h"  // the shared ops (tl::gpu) resolve to cuda here
+
+#include <cstdio>
+#include <vector>
+
+namespace cu = tl::cuda;
+using kop = tl::gpu::kop;
+
+namespace {
+
+struct buf {
+  void* native = nullptr;
+  float* host = nullptr;
+  int64_t bytes = 0;
+  explicit buf(int64_t b) : bytes(b) { native = cu::alloc(b, &host); }
+  ~buf() { cu::release(native, bytes, host); }
+  buf(const buf&) = delete;
+  operator void*() const { return native; }
+};
+
+constexpr int64_t kOff = 64;  // bytes; 16B-aligned so the tiled GEMM still takes it
+
+void elementwise() {
+  const int64_t n = 1000;
+  buf a(n * 4 + kOff), b(n * 4 + kOff), o(n * 4 + kOff);
+  for (kop op : {kop::add, kop::sub, kop::mul, kop::div, kop::pow_}) {
+    tl::gpu::binary(op, {a, kOff}, {b, 0}, {o, kOff}, n, 2.0f, 1.0f);
+  }
+  for (kop op : {kop::badd, kop::bsub, kop::bmul, kop::bdiv, kop::bpow}) {
+    tl::gpu::binary_bcast(op, {a, kOff}, 40, 1, {b, 0}, 0, 1, {o, kOff}, 25, 40,
+                          1.0f, 0.0f);
+  }
+  const int64_t shape[3] = {5, 8, 25}, as[3] = {200, 25, 1}, bs[3] = {0, 25, 1};
+  for (kop op : {kop::badd, kop::bsub, kop::bmul, kop::bdiv, kop::bpow}) {
+    tl::gpu::binary_bcast_nd(op, {a, kOff}, as, {b, 0}, bs, {o, kOff}, shape, 3, n, 1.0f, 0.0f);
+  }
+  using cu::cmp_op;
+  for (cmp_op op : {cmp_op::gt, cmp_op::lt, cmp_op::ge, cmp_op::le, cmp_op::eq,
+                    cmp_op::ne}) {
+    tl::gpu::compare(op, {a, kOff}, {b, 0}, {o, kOff}, n, 1);
+  }
+}
+
+void gemm_bias() {
+  struct shape { int64_t m, n, k; };
+  for (shape s : {shape{8, 8, 8}, {64, 64, 64}, {256, 512, 512},
+                  {256, 2048, 512}, {512, 896, 896}}) {
+    buf a(s.m * s.k * 4 + kOff), b(s.k * s.n * 4 + kOff), bias(s.n * 4),
+        o(s.m * s.n * 4 + kOff);
+    for (int layout = 0; layout < 4; layout++) {
+      const bool ta = layout & 1, tb = layout & 2;
+      tl::gpu::gemm_bias({a, kOff}, ta ? s.m : s.k, ta, {b, kOff}, tb ? s.k : s.n, tb, {bias, 0}, {o, kOff}, s.m, s.n, s.k, 1.0f, 0.0f);
+    }
+  }
+}
+
+// The four ops that zero their output before scattering into it.
+void zero_then_scatter() {
+  const int64_t a_shape[2] = {4, 6}, out_shape[2] = {4, 10};
+  buf a(24 * 4 + kOff), o(40 * 4 + kOff);
+  tl::gpu::pad({a, kOff}, {o, kOff}, a_shape, out_shape, 2, 1, 2, 24, 40);
+
+  const int64_t w_shape[3] = {4, 4, 3}, f_shape[2] = {4, 6};
+  buf w(48 * 4 + kOff), f(24 * 4 + kOff);
+  tl::gpu::fold({w, kOff}, {f, kOff}, w_shape, f_shape, 3, 1, 1, 48, 24);
+
+  buf idx(8 * 4 + kOff), vals(8 * 16 * 4 + kOff), table(32 * 16 * 4 + kOff);
+  tl::gpu::index_add({idx, kOff}, {vals, kOff}, {table, kOff}, 16, 8, 32 * 16);
+  tl::gpu::scatter_to_axis({idx, kOff}, {vals, kOff}, {table, kOff}, 8, 16);
+}
+
+void llm() {
+  for (int64_t m : {8, 96, 512}) {
+    for (int64_t n : {896, 4864}) {
+      const int64_t k = 896;
+      buf a(m * k * 4), B(n * k * 2), o(m * n * 4);
+      tl::gpu::gemm_bf16_nt({a, 0}, {B, 0}, {o, 0}, m, n, k);
+    }
+  }
+  for (int64_t D : {64, 128}) {
+    const int64_t hq = 14, hkv = 2, kv_max = 4096;
+    buf q(hq * D * 4), K(hkv * kv_max * D * 2), V(hkv * kv_max * D * 2),
+        o(hq * D * 4);
+    for (int64_t ctx : {17, 700, 4000}) {
+      tl::gpu::attn_decode({q, 0}, {K, 0}, {V, 0}, {o, 0}, hq, hkv, ctx, kv_max, D, 0.125f, true);
+    }
+  }
+}
+
+// The graph-capture group: device-position variants, recorded then replayed.
+void capture(int64_t D) {
+  const int64_t hq = 14, hkv = 2, kv_max = 2048;
+  buf pos(4), x(hq * D * 4), xo(hq * D * 4), knew(hkv * D * 4), vnew(hkv * D * 4),
+      K(hkv * kv_max * D * 4), V(hkv * kv_max * D * 4), o(hq * D * 4),
+      partials(cu::attn_dpos_partials_bytes(hq, kv_max, D));
+  cu::upload_u32(pos, 5);
+  if (!cu::capture_begin()) return;
+  tl::gpu::rope_dpos({x, 0}, {xo, 0}, hq, 1, D, {pos, 0}, 10000.0f);
+  tl::gpu::kv_append_dpos({K, 0}, {V, 0}, {knew, 0}, {vnew, 0}, {pos, 0}, kv_max, hkv, D);
+  tl::gpu::attn_decode_dpos({xo, 0}, {K, 0}, {V, 0}, {o, 0}, hq, hkv, {pos, 0}, kv_max, D, 0.125f, {partials, 0});
+  tl::gpu::incr_u32({pos, 0});
+  auto graph = cu::capture_end();
+  cu::graph_launch(graph);
+  cu::flush();
+  cu::graph_destroy(graph);
+}
+
+}  // namespace
+
+int main() {
+  if (!cu::available()) return std::printf("no CUDA driver: nothing to trace\n"), 0;
+  elementwise();
+  gemm_bias();
+  zero_then_scatter();
+  llm();
+  capture(64);
+  capture(128);
+  cu::flush();
+  return 0;
+}

bench/cuda/speed/bench_attn_bwd.cpp

@@ -12,7 +12,7 @@
 #ifndef TENSORLIB_CUDA
 #define TENSORLIB_CUDA
 #endif
-#include "cuda.h"
+#include "gpu.h"  // cuda.h plus the shared ops (tl::gpu resolves to cuda here)
 
 #include <algorithm>
 #include <chrono>
@@ -72,13 +72,13 @@ int main() {
     fill_random(hG, n, 4);
 
     auto run_fwd = [&] {
-      attn_prefill(q, K, V, out, s.H, s.H, s.T, s.T, s.D, scale);
+      tl::gpu::attn_prefill({q, 0}, {K, 0}, {V, 0}, {out, 0}, s.H, s.H, s.T, s.T, s.D, scale);
     };
     auto run_dq = [&] {
-      attn_prefill_dq(q, K, V, dO, out, dq, stats, s.H, s.T, s.D, scale);
+      tl::gpu::attn_prefill_dq({q, 0}, {K, 0}, {V, 0}, {dO, 0}, {out, 0}, {dq, 0}, {stats, 0}, s.H, s.T, s.D, scale);
     };
     auto run_dkv = [&] {
-      attn_prefill_dkv(q, K, V, dO, stats, dK, dV, s.H, s.T, s.D, scale);
+      tl::gpu::attn_prefill_dkv({q, 0}, {K, 0}, {V, 0}, {dO, 0}, {stats, 0}, {dK, 0}, {dV, 0}, s.H, s.T, s.D, scale);
     };
     auto time_it = [&](auto&& fn) {
       fn();