tsotchke · tsotchke · Aug 30, 2026 · Aug 30, 2026 · Aug 30, 2026 · Aug 30, 2026
diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml
@@ -34,6 +34,12 @@ jobs:
       - name: Documentation links
         run: python3 scripts/check_docs_links.py
 
+      - name: tc-cuda v1 subset authority
+        run: |
+          python3 scripts/check_tc_cuda_subset.py
+          python3 scripts/check_tc_cuda_subset_selftest.py
+          python3 scripts/tc_cuda_selftest.py
+
       - name: Public release privacy
         run: |
           python3 scripts/check_release_privacy_selftest.py

diff --git a/docs/README.md b/docs/README.md
@@ -57,6 +57,9 @@ are the entry points; everything below goes deeper.
   per-dtype hardware gates, SDK gates, and how the dispatch picks a path.
 - **[cuda_comparison.md](cuda_comparison.md)** — direct
   cuBLAS / cuDNN / CUTLASS / NCCL / Triton ↔ tensorcore equivalents.
+- **[tc-cuda/README.md](tc-cuda/README.md)** — the tc-cuda v1 CUDA subset
+  authority: the 55-construct accept-list, the unsupported list, and the
+  fail-closed policy, all generated from `docs/tc-cuda/subset.v1.json`.
 
 ### Kernels
 

diff --git a/docs/ci_and_scripts.md b/docs/ci_and_scripts.md
@@ -27,6 +27,12 @@ on macOS runners. Core gates include:
 5. **`scripts/ci_python_smoke.sh`** — sets up a venv, installs the binding
    editable, asserts `tc.version()`, the diagnostic helpers, and the
    tensorops kernel selector.
+6. **`scripts/check_tc_cuda_subset.py`** + **`scripts/check_tc_cuda_subset_selftest.py`**
+   — the tc-cuda v1 subset authority gate. The validator loads
+   `docs/tc-cuda/subset.v1.json` (the single source of truth for the accept-list)
+   and fails closed on duplicate ids, unknown categories, any status other than
+   `supported`/`unsupported`, or any construct that would be silently accepted.
+   The selftest proves the validator itself rejects those malformed manifests.
 
 This is the gate. PRs need it green to merge.
 

diff --git a/docs/tc-cuda/README.md b/docs/tc-cuda/README.md
@@ -0,0 +1,68 @@
+# tc-cuda v1 CUDA subset
+
+This directory holds the machine-readable authority for the tc-cuda v1 CUDA
+subset. The single source of truth is:
+
+**`docs/tc-cuda/subset.v1.json`**
+
+The tc-cuda compiler's accept-list and all subset documentation are generated
+from that one artefact. Doc and code cannot drift because they are one
+artefact. This README is a summary only and never overrides the JSON.
+
+## Scope
+
+tc-cuda accepts CUDA kernel source and CUDA host API calls and executes them on
+tensorcore's backends. Its default behaviour on a construct it cannot handle is
+to **refuse to build** — never to emit a binary whose kernel silently does
+nothing. See the design document for the full requirements:
+
+- `docs/design/tc-cuda-universal-substrate-20260827.md`
+
+## Supported constructs — 55
+
+The subset is the union of what `lib/cuda/*.cu` and the Kimi engine's `*.cu`
+actually use, surveyed not guessed. Category counts:
+
+| Category | Count |
+| --- | ---: |
+| A. Function and declaration forms | 4 |
+| B. Launch geometry and indexing | 3 |
+| C. Memory | 4 |
+| D. Synchronisation and warp collectives | 4 |
+| E. Atomics | 1 |
+| F. Half precision | 8 |
+| G. fp32 libdevice math | 13 |
+| H. Rounding-explicit and fast intrinsics | 4 |
+| I. Integer and bit arithmetic | 2 |
+| J. Host runtime API families | 12 |
+| **Total** | **55** |
+
+Each supported entry in the JSON carries a stable `id`, the CUDA `name`,
+`status: "supported"`, and concise `semantics`.
+
+## Unsupported constructs
+
+Every construct in the design's "explicitly out of scope in v1" list is encoded
+in the JSON with `status: "unsupported"` and a stable `id`. Each is a **build
+error that names the construct** in CUDA's own vocabulary with file:line. This
+includes the performance hints `__ldg` and `__launch_bounds__`, which are
+deliberately rejected rather than silently ignored.
+
+## Fail-closed policy
+
+- **Single authority.** `subset.v1.json` is the only source of truth; the
+  compiler accept-list is generated from it.
+- **Fail closed.** Any construct not listed as `supported` is a build error
+  naming the construct. There is no best-effort mode and no silent fallback.
+- **Uniqueness.** Every `id` is unique across supported and unsupported.
+  Duplicate ids, unknown categories, or any status other than
+  `supported`/`unsupported` are validation failures.
+- **No silent acceptance.** A construct that is present in source but absent
+  from the supported list is never accepted; it is named and rejected.
+
+## Source checker
+
+Run `python3 scripts/tc_cuda.py check SOURCE.cu --manifest-output manifest.json`.
+The emitted `checked` status proves source validation only. CUDA-to-Metal
+lowering and execution are not claimed until their separate conformance gates
+pass.
diff --git a/docs/tc-cuda/subset.v1.json b/docs/tc-cuda/subset.v1.json
@@ -0,0 +1,107 @@
+{
+  "schema": "tensorcore.tc-cuda.subset.v1",
+  "version": "1.0.0",
+  "design_ref": "docs/design/tc-cuda-universal-substrate-20260827.md#3",
+  "total": 55,
+  "category_counts": {
+    "A": 4,
+    "B": 3,
+    "C": 4,
+    "D": 4,
+    "E": 1,
+    "F": 8,
+    "G": 13,
+    "H": 4,
+    "I": 2,
+    "J": 12
+  },
+  "policy": {
+    "authority": "This file is the sole machine-readable authority for the tc-cuda v1 subset. The compiler accept-list and all documentation are generated from it; doc and code cannot drift because they are one artefact.",
+    "fail_closed": "Any construct not listed here with status=supported is a build error that names the construct in CUDA's own vocabulary with file:line. There is no best-effort mode and no silent fallback.",
+    "uniqueness": "Every id is unique across supported and unsupported. Duplicate ids, unknown categories, or a status other than supported|unsupported are validation failures."
+  },
+  "supported": [
+    { "id": "A1", "name": "__global__", "status": "supported", "semantics": "Kernel entry point: __global__ void f(...). Enumerated in the closed kernel manifest; link fails if any is not lowered (R1)." },
+    { "id": "A2", "name": "__device__", "status": "supported", "semantics": "Device helper function, including inline and static linkage." },
+    { "id": "A3", "name": "template<int D>", "status": "supported", "semantics": "Non-type integral template parameter; explicit instantiation only. Type-parametric templates are out of scope." },
+    { "id": "A4", "name": "__restrict__", "status": "supported", "semantics": "Restriction qualifier on pointer parameters; must lower to a real noalias, not be dropped." },
+    { "id": "B1", "name": "threadIdx/blockIdx/blockDim/gridDim", "status": "supported", "semantics": "Launch geometry indexing with .x .y .z components; .y is used by training.cu and kimi_moe_cuda.cu." },
+    { "id": "B2", "name": "<<<grid, block>>>", "status": "supported", "semantics": "Kernel launch configuration with grid and block dimensions." },
+    { "id": "B3", "name": "<<<grid, block, dynamic_smem_bytes>>>", "status": "supported", "semantics": "Launch with dynamic shared-memory byte count; paired with C2." },
+    { "id": "C1", "name": "__shared__", "status": "supported", "semantics": "Static-extent shared-memory array." },
+    { "id": "C2", "name": "extern __shared__", "status": "supported", "semantics": "Dynamic shared-memory array; sized via the third launch argument (B3)." },
+    { "id": "C3", "name": "__device__ variable + cudaMemcpyToSymbol/FromSymbol", "status": "supported", "semantics": "Module-scope device variable with symbol copy in both directions." },
+    { "id": "C4", "name": "global pointer load/store", "status": "supported", "semantics": "Global memory load/store, including const T* __restrict__." },
+    { "id": "D1", "name": "__syncthreads()", "status": "supported", "semantics": "Block-level barrier synchronisation." },
+    { "id": "D2", "name": "__shfl_sync", "status": "supported", "semantics": "Warp shuffle with explicit member mask." },
+    { "id": "D3", "name": "__shfl_xor_sync", "status": "supported", "semantics": "Warp shuffle-xor; highest-risk warp collective, see design §7." },
+    { "id": "D4", "name": "__shfl_down_sync", "status": "supported", "semantics": "Warp shuffle-down with explicit member mask." },
+    { "id": "E1", "name": "atomicAdd(float*, float)", "status": "supported", "semantics": "Global-memory float atomic add. Shared-memory and integer atomics are out of scope." },
+    { "id": "F1", "name": "__half", "status": "supported", "semantics": "Half-precision scalar, host and device storage; host/device attribute parity is part of the contract (R8)." },
+    { "id": "F2", "name": "__half2", "status": "supported", "semantics": "Packed half-precision pair." },
+    { "id": "F3", "name": "__half2float / __float2half", "status": "supported", "semantics": "Half/float conversion, both __host__ __device__." },
+    { "id": "F4", "name": "__float2half_rn", "status": "supported", "semantics": "Round-to-nearest float to half." },
+    { "id": "F5", "name": "__float2half2_rn", "status": "supported", "semantics": "Round-to-nearest float to packed half2." },
+    { "id": "F6", "name": "__hfma2", "status": "supported", "semantics": "Packed half fused multiply-add." },
+    { "id": "F7", "name": "__hmul2", "status": "supported", "semantics": "Packed half multiply." },
+    { "id": "F8", "name": "__ushort_as_half", "status": "supported", "semantics": "Bit-cast between unsigned short and half; must not round-trip through float." },
+    { "id": "G1", "name": "sqrtf", "status": "supported", "semantics": "fp32 square root; needs its own numerical golden." },
+    { "id": "G2", "name": "rsqrtf", "status": "supported", "semantics": "fp32 reciprocal square root; needs its own numerical golden." },
+    { "id": "G3", "name": "expf", "status": "supported", "semantics": "fp32 natural exponential; needs its own numerical golden." },
+    { "id": "G4", "name": "logf", "status": "supported", "semantics": "fp32 natural logarithm; needs its own numerical golden." },
+    { "id": "G5", "name": "tanhf", "status": "supported", "semantics": "fp32 hyperbolic tangent; needs its own numerical golden." },
+    { "id": "G6", "name": "powf", "status": "supported", "semantics": "fp32 power; needs its own numerical golden." },
+    { "id": "G7", "name": "sinf", "status": "supported", "semantics": "fp32 sine; needs its own numerical golden." },
+    { "id": "G8", "name": "cosf", "status": "supported", "semantics": "fp32 cosine; needs its own numerical golden." },
+    { "id": "G9", "name": "fmaxf", "status": "supported", "semantics": "fp32 maximum; needs its own numerical golden." },
+    { "id": "G10", "name": "fminf", "status": "supported", "semantics": "fp32 minimum; needs its own numerical golden." },
+    { "id": "G11", "name": "fabsf", "status": "supported", "semantics": "fp32 absolute value; needs its own numerical golden." },
+    { "id": "G12", "name": "floorf", "status": "supported", "semantics": "fp32 floor; needs its own numerical golden." },
+    { "id": "G13", "name": "fmaf", "status": "supported", "semantics": "fp32 fused multiply-add; contraction policy is declared per translation unit (R4)." },
+    { "id": "H1", "name": "__expf", "status": "supported", "semantics": "Fast exponential; maps to metal::fast::exp, not metal::exp." },
+    { "id": "H2", "name": "__logf", "status": "supported", "semantics": "Fast logarithm." },
+    { "id": "H3", "name": "__fmul_rn", "status": "supported", "semantics": "Round-to-nearest multiply; contraction barrier — must not be fused (R4)." },
+    { "id": "H4", "name": "__fadd_rn", "status": "supported", "semantics": "Round-to-nearest add; contraction barrier — must not be fused (R4)." },
+    { "id": "I1", "name": "integer/bit arithmetic", "status": "supported", "semantics": "uint8_t / unsigned short / int / long / size_t arithmetic, shifts, masks, casts; the whole of dequant_q4_t_kernel." },
+    { "id": "I2", "name": "integer division and modulo", "status": "supported", "semantics": "Integer division and modulo by a runtime value, e.g. idx / nblocks, idx % nblocks." },
+    { "id": "J1", "name": "cudaMalloc", "status": "supported", "semantics": "Device memory allocation." },
+    { "id": "J2", "name": "cudaMallocHost", "status": "supported", "semantics": "Pinned host memory allocation." },
+    { "id": "J3", "name": "cudaMallocManaged", "status": "supported", "semantics": "Unified managed memory; nearly free on Apple unified memory." },
+    { "id": "J4", "name": "cudaFree", "status": "supported", "semantics": "Device memory deallocation." },
+    { "id": "J5", "name": "cudaMemcpy / cudaMemcpyAsync", "status": "supported", "semantics": "Host/device memory copy in all 4 directions, sync and async." },
+    { "id": "J6", "name": "cudaMemset", "status": "supported", "semantics": "Device memory fill." },
+    { "id": "J7", "name": "cudaMemcpyToSymbol / cudaMemcpyFromSymbol", "status": "supported", "semantics": "Symbol-based device variable copy, both directions." },
+    { "id": "J8", "name": "cudaStream*", "status": "supported", "semantics": "Stream Create/WithFlags/Destroy/Synchronize." },
+    { "id": "J9", "name": "cudaEvent*", "status": "supported", "semantics": "Event Create/Record/Synchronize/ElapsedTime/Destroy." },
+    { "id": "J10", "name": "cudaDeviceSynchronize", "status": "supported", "semantics": "Device-wide synchronisation; keeps returning tcCudaErrorKernelNotLowered if a kernel was not lowered (R3)." },
+    { "id": "J11", "name": "error surface", "status": "supported", "semantics": "cudaGetLastError, cudaGetErrorName, cudaSuccess, cudaError_t; no path from no-kernel to success (R3)." },
+    { "id": "J12", "name": "device query", "status": "supported", "semantics": "cudaGetDevice/SetDevice/GetDeviceCount/GetDeviceProperties/cudaPointerGetAttributes/cudaFuncSetAttribute." }
+  ],
+  "unsupported": [
+    { "id": "U1", "name": "wmma / mma_sync / tensor-core intrinsics", "status": "unsupported", "reason": "No backend mapping in v1; build error naming the construct." },
+    { "id": "U2", "name": "inline PTX (asm volatile)", "status": "unsupported", "reason": "Source-only route; PTX is not ingested (non-goal)." },
+    { "id": "U3", "name": "cp.async", "status": "unsupported", "reason": "Async copy intrinsic has no v1 mapping." },
+    { "id": "U4", "name": "ldmatrix", "status": "unsupported", "reason": "Matrix load intrinsic has no v1 mapping." },
+    { "id": "U5", "name": "cooperative groups", "status": "unsupported", "reason": "No v1 lowering." },
+    { "id": "U6", "name": "CUDA graphs", "status": "unsupported", "reason": "No v1 lowering." },
+    { "id": "U7", "name": "dynamic parallelism", "status": "unsupported", "reason": "No v1 lowering." },
+    { "id": "U8", "name": "texture and surface objects", "status": "unsupported", "reason": "No v1 lowering." },
+    { "id": "U9", "name": "__constant__", "status": "unsupported", "reason": "Constant memory space not in v1." },
+    { "id": "U10", "name": "warp-vote (__ballot_sync, __any_sync, __all_sync)", "status": "unsupported", "reason": "No v1 mapping; silent_noop.cu asserts rejection." },
+    { "id": "U11", "name": "__syncwarp", "status": "unsupported", "reason": "Warp-level barrier not in v1." },
+    { "id": "U12", "name": "__ldg", "status": "unsupported", "reason": "Performance hint; rejecting avoids a silent performance cliff (R1)." },
+    { "id": "U13", "name": "shared-memory atomics", "status": "unsupported", "reason": "Only global float atomicAdd (E1) is in scope." },
+    { "id": "U14", "name": "integer atomics", "status": "unsupported", "reason": "Only global float atomicAdd (E1) is in scope." },
+    { "id": "U15", "name": "atomicCAS / atomicExch / atomicMax", "status": "unsupported", "reason": "Only global float atomicAdd (E1) is in scope." },
+    { "id": "U16", "name": "double-precision math", "status": "unsupported", "reason": "No v1 fp64 device path." },
+    { "id": "U17", "name": "device-side printf", "status": "unsupported", "reason": "No v1 device I/O." },
+    { "id": "U18", "name": "device-side malloc/free/assert", "status": "unsupported", "reason": "No v1 device heap or assert." },
+    { "id": "U19", "name": "recursion", "status": "unsupported", "reason": "Device recursion not lowered in v1." },
+    { "id": "U20", "name": "virtual functions and RTTI in device code", "status": "unsupported", "reason": "No v1 device vtable/RTTI." },
+    { "id": "U21", "name": "type-parametric templates", "status": "unsupported", "reason": "Only non-type integral templates (A3) are in scope." },
+    { "id": "U22", "name": "__launch_bounds__", "status": "unsupported", "reason": "Performance hint; rejecting avoids a silent performance cliff (R1)." },
+    { "id": "U23", "name": "multi-GPU peer access", "status": "unsupported", "reason": "No v1 peer mapping." },
+    { "id": "U24", "name": "stream callbacks", "status": "unsupported", "reason": "No v1 callback host hook." },
+    { "id": "U25", "name": "float4 / double2 vector types", "status": "unsupported", "reason": "Vector types not in v1." }
+  ]
+}