diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index c1507c7..55a631b 100644 --- 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 index 66dfdd8..dab62b3 100644 --- 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 index dcdee7d..3e520e1 100644 --- 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 new file mode 100644 index 0000000..44fb727 --- /dev/null +++ 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 new file mode 100644 index 0000000..cf46eef --- /dev/null +++ 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", "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": "<<>>", "status": "supported", "semantics": "Kernel launch configuration with grid and block dimensions." }, + { "id": "B3", "name": "<<>>", "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." } + ] +} diff --git a/scripts/check_tc_cuda_subset.py b/scripts/check_tc_cuda_subset.py new file mode 100644 index 0000000..0f6f86a --- /dev/null +++ b/scripts/check_tc_cuda_subset.py @@ -0,0 +1,275 @@ +#!/usr/bin/env python3 +"""Fail-closed validator for the tc-cuda v1 CUDA subset authority. + +`docs/tc-cuda/subset.v1.json` is the sole machine-readable authority for the +tc-cuda v1 subset. The compiler accept-list and all subset documentation are +generated from it, so this validator is the gate that keeps that single +artifact self-consistent. It is fail-closed: any duplicate id, unknown +category, wrong status, missing required string, or unknown key is a +validation failure (exit 1), never a silent acceptance. + +Usage: + python3 scripts/check_tc_cuda_subset.py [path-to-subset.json] + +Defaults to docs/tc-cuda/subset.v1.json relative to the repository root. +""" + +from __future__ import annotations + +import argparse +import json +import pathlib +import sys +from typing import Any + + +ROOT = pathlib.Path(__file__).resolve().parents[1] +DEFAULT_MANIFEST = ROOT / "docs" / "tc-cuda" / "subset.v1.json" + +SCHEMA = "tensorcore.tc-cuda.subset.v1" +VERSION = "1.0.0" + +# Top-level keys the manifest is allowed to carry. Anything else is an +# unknown key and a validation failure (fail-closed). +ALLOWED_TOP_LEVEL = { + "schema", + "version", + "design_ref", + "total", + "category_counts", + "policy", + "supported", + "unsupported", +} + +# Per-entry keys. Supported entries carry `semantics`; unsupported entries +# carry `reason`. Both carry the common id/name/status. +COMMON_ENTRY_KEYS = {"id", "name", "status"} +ALLOWED_SUPPORTED_KEYS = COMMON_ENTRY_KEYS | {"semantics"} +ALLOWED_UNSUPPORTED_KEYS = COMMON_ENTRY_KEYS | {"reason"} + +VALID_STATUS = {"supported", "unsupported"} + +# Supported categories are A-J (function forms, launch geometry, memory, +# synchronisation, atomics, half, fp32 math, intrinsics, integer, host API). +# Unsupported entries use the U prefix. +SUPPORTED_CATEGORY_LETTERS = set("ABCDEFGHIJ") +UNSUPPORTED_CATEGORY_LETTER = "U" + + +def parse_args() -> argparse.Namespace: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument( + "manifest", + nargs="?", + type=pathlib.Path, + default=DEFAULT_MANIFEST, + help="path to subset.v1.json (default: docs/tc-cuda/subset.v1.json)", + ) + return parser.parse_args() + + +def load_manifest(path: pathlib.Path) -> Any: + try: + text = path.read_text(encoding="utf-8") + except OSError as exc: + raise SystemExit(f"could not read subset manifest {path}: {exc}") from exc + try: + return json.loads(text) + except json.JSONDecodeError as exc: + raise SystemExit(f"subset manifest is not valid JSON: {exc}") from exc + + +def fail(errors: list[str]) -> int: + print("tc-cuda subset manifest invalid:", file=sys.stderr) + for error in errors: + print(f" - {error}", file=sys.stderr) + return 1 + + +def is_nonempty_str(value: Any) -> bool: + return isinstance(value, str) and value.strip() != "" + + +def is_positive_int(value: Any) -> bool: + return isinstance(value, int) and not isinstance(value, bool) and value >= 1 + + +def check_top_level(errors: list[str], data: dict[str, Any]) -> None: + unknown = sorted(set(data) - ALLOWED_TOP_LEVEL) + if unknown: + errors.append(f"unknown top-level key(s): {unknown!r}") + + if data.get("schema") != SCHEMA: + errors.append(f"schema must be {SCHEMA!r}, got {data.get('schema')!r}") + if data.get("version") != VERSION: + errors.append(f"version must be {VERSION!r}, got {data.get('version')!r}") + if not is_nonempty_str(data.get("design_ref")): + errors.append("design_ref must be a non-empty string") + + if not is_positive_int(data.get("total")): + errors.append(f"total must be a positive integer, got {data.get('total')!r}") + + if not isinstance(data.get("supported"), list): + errors.append("supported must be a list") + if not isinstance(data.get("unsupported"), list): + errors.append("unsupported must be a list") + + +def check_category_counts(errors: list[str], data: dict[str, Any]) -> dict[str, int]: + counts = data.get("category_counts") + if not isinstance(counts, dict): + errors.append("category_counts must be an object") + return {} + total = 0 + for key, value in counts.items(): + if not is_nonempty_str(key): + errors.append(f"category_counts key must be a non-empty string, got {key!r}") + continue + if not is_positive_int(value): + errors.append( + f"category_counts[{key!r}] must be a positive integer, got {value!r}" + ) + continue + total += value + if is_positive_int(data.get("total")) and total != data["total"]: + errors.append( + f"category_counts sum {total} must equal total {data['total']}" + ) + return counts if isinstance(counts, dict) else {} + + +def check_entry( + errors: list[str], + entry: Any, + index: int, + section: str, + allowed_keys: set[str], + extra_required: tuple[str, ...], + expected_status: str, +) -> str | None: + """Validate one entry; return its id if usable, else None.""" + label = f"{section}[{index}]" + if not isinstance(entry, dict): + errors.append(f"{label} must be an object, got {type(entry).__name__}") + return None + + unknown = sorted(set(entry) - allowed_keys) + if unknown: + errors.append(f"{label} has unknown key(s): {unknown!r}") + + entry_id = entry.get("id") + if not is_nonempty_str(entry_id): + errors.append(f"{label}.id must be a non-empty string, got {entry_id!r}") + entry_id = None + + if not is_nonempty_str(entry.get("name")): + errors.append(f"{label}.name must be a non-empty string") + + if entry.get("status") != expected_status: + errors.append( + f"{label}.status must be {expected_status!r}, got {entry.get('status')!r}" + ) + + for key in extra_required: + if not is_nonempty_str(entry.get(key)): + errors.append(f"{label}.{key} must be a non-empty string") + + return entry_id + + +def check_supported( + errors: list[str], + data: dict[str, Any], + counts: dict[str, int], +) -> list[str]: + entries = data.get("supported") + if not isinstance(entries, list): + return [] + ids: list[str] = [] + per_category: dict[str, int] = {} + for index, entry in enumerate(entries): + entry_id = check_entry( + errors, entry, index, "supported", + ALLOWED_SUPPORTED_KEYS, ("semantics",), "supported", + ) + if entry_id is None: + continue + ids.append(entry_id) + letter = entry_id[0] + if letter not in SUPPORTED_CATEGORY_LETTERS: + errors.append( + f"supported[{index}].id {entry_id!r} must start with one of " + f"{sorted(SUPPORTED_CATEGORY_LETTERS)!r}" + ) + else: + if letter not in counts: + errors.append( + f"supported[{index}].id {entry_id!r} category {letter!r} " + "is not declared in category_counts" + ) + per_category[letter] = per_category.get(letter, 0) + 1 + + for letter, seen in sorted(per_category.items()): + declared = counts.get(letter) + if isinstance(declared, int) and declared != seen: + errors.append( + f"category_counts[{letter!r}] is {declared} but {seen} supported " + f"entries use category {letter!r}" + ) + return ids + + +def check_unsupported(errors: list[str], data: dict[str, Any]) -> list[str]: + entries = data.get("unsupported") + if not isinstance(entries, list): + return [] + ids: list[str] = [] + for index, entry in enumerate(entries): + entry_id = check_entry( + errors, entry, index, "unsupported", + ALLOWED_UNSUPPORTED_KEYS, ("reason",), "unsupported", + ) + if entry_id is None: + continue + ids.append(entry_id) + if entry_id[0] != UNSUPPORTED_CATEGORY_LETTER: + errors.append( + f"unsupported[{index}].id {entry_id!r} must start with " + f"{UNSUPPORTED_CATEGORY_LETTER!r}" + ) + return ids + + +def check_duplicates(errors: list[str], supported_ids: list[str], unsupported_ids: list[str]) -> None: + all_ids = supported_ids + unsupported_ids + seen: dict[str, str] = {} + for entry_id in all_ids: + if entry_id in seen: + errors.append(f"duplicate id {entry_id!r} (first seen in {seen[entry_id]!r})") + else: + seen[entry_id] = "supported" if entry_id in supported_ids else "unsupported" + + +def main() -> int: + args = parse_args() + data = load_manifest(args.manifest) + if not isinstance(data, dict): + return fail(["top-level manifest must be a JSON object"]) + + errors: list[str] = [] + check_top_level(errors, data) + counts = check_category_counts(errors, data) + supported_ids = check_supported(errors, data, counts) + unsupported_ids = check_unsupported(errors, data) + check_duplicates(errors, supported_ids, unsupported_ids) + + if errors: + return fail(errors) + + print(f"tc-cuda subset manifest OK: {args.manifest} ({data['total']} entries)") + return 0 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/scripts/check_tc_cuda_subset_selftest.py b/scripts/check_tc_cuda_subset_selftest.py new file mode 100644 index 0000000..a5fc4a4 --- /dev/null +++ b/scripts/check_tc_cuda_subset_selftest.py @@ -0,0 +1,193 @@ +#!/usr/bin/env python3 +"""Regression selftest for the tc-cuda v1 subset validator. + +`scripts/check_tc_cuda_subset.py` is the fail-closed gate that keeps +`docs/tc-cuda/subset.v1.json` self-consistent. This selftest proves two things: + + 1. The committed real manifest passes the validator (exit 0). + 2. Each representative fail-closed mutation is REJECTED (exit 1) — i.e. the + validator never silently accepts a duplicate, unknown, or mis-declared + construct. + +It uses only the Python stdlib and writes mutated copies into a temporary +directory (the committed manifest is never modified). It exits 0 only when the +real manifest passes and every mutation is rejected. + +Usage: + python3 scripts/check_tc_cuda_subset_selftest.py +""" + +from __future__ import annotations + +import copy +import importlib.util +import json +import pathlib +import subprocess +import sys +import tempfile + + +ROOT = pathlib.Path(__file__).resolve().parents[1] +VALIDATOR = ROOT / "scripts" / "check_tc_cuda_subset.py" +MANIFEST = ROOT / "docs" / "tc-cuda" / "subset.v1.json" + + +def load_validator_module(): + spec = importlib.util.spec_from_file_location("check_tc_cuda_subset", VALIDATOR) + module = importlib.util.module_from_spec(spec) + spec.loader.exec_module(module) + return module + + +def load_real_manifest() -> dict: + return json.loads(MANIFEST.read_text(encoding="utf-8")) + + +def run_validator(path: pathlib.Path) -> int: + proc = subprocess.run( + [sys.executable, str(VALIDATOR), str(path)], + capture_output=True, + text=True, + ) + return proc.returncode + + +def first_supported(data: dict) -> dict: + return data["supported"][0] + + +def first_unsupported(data: dict) -> dict: + return data["unsupported"][0] + + +def build_mutations(base: dict) -> list[tuple[str, dict]]: + """Return (label, mutated-manifest) pairs, each exercising one fail-closed rule.""" + mutations: list[tuple[str, dict]] = [] + + # 1. Unknown top-level key. + m = copy.deepcopy(base) + m["bogus_top_level"] = "should be rejected" + mutations.append(("unknown top-level key", m)) + + # 2. Wrong schema. + m = copy.deepcopy(base) + m["schema"] = "tensorcore.tc-cuda.subset.v2" + mutations.append(("wrong schema", m)) + + # 3. Wrong version. + m = copy.deepcopy(base) + m["version"] = "9.9.9" + mutations.append(("wrong version", m)) + + # 4. Incorrect total (mismatched with category_counts sum). + m = copy.deepcopy(base) + m["total"] = base["total"] + 1 + mutations.append(("incorrect total", m)) + + # 5. Incorrect category_counts (a declared count no longer matches entries). + m = copy.deepcopy(base) + m["category_counts"] = copy.deepcopy(base["category_counts"]) + m["category_counts"]["A"] = base["category_counts"]["A"] + 1 + mutations.append(("incorrect category_counts", m)) + + # 6. Duplicate supported id (a second entry reuses the first supported id). + m = copy.deepcopy(base) + dup = copy.deepcopy(first_supported(m)) + m["supported"].append(dup) + mutations.append(("duplicate supported id", m)) + + # 7. Supported/unsupported overlap (an unsupported entry reuses a supported id). + m = copy.deepcopy(base) + overlap = { + "id": first_supported(m)["id"], + "name": "cross-section duplicate", + "status": "unsupported", + "reason": "should be rejected as a duplicate id", + } + m["unsupported"].append(overlap) + mutations.append(("supported/unsupported id overlap", m)) + + # 8. Invalid status on a supported entry. + m = copy.deepcopy(base) + first_supported(m)["status"] = "maybe" + mutations.append(("invalid supported status", m)) + + # 9. Unknown entry key on a supported entry. + m = copy.deepcopy(base) + first_supported(m)["extra_key"] = "should be rejected" + mutations.append(("unknown supported entry key", m)) + + # 10. Missing required non-string value (name not a string). + m = copy.deepcopy(base) + first_supported(m)["name"] = 12345 + mutations.append(("non-string supported name", m)) + + # 11. Missing required string value (semantics absent). + m = copy.deepcopy(base) + del first_supported(m)["semantics"] + mutations.append(("missing supported semantics", m)) + + # 12. Unknown category letter on a supported entry. + m = copy.deepcopy(base) + first_supported(m)["id"] = "Z1" + mutations.append(("unknown supported category", m)) + + # 13. Category membership mismatch (supported id in a category not in category_counts). + m = copy.deepcopy(base) + m["category_counts"] = {k: v for k, v in base["category_counts"].items() if k != "A"} + # total must still be a positive int so per-category comparison is exercised. + mutations.append(("category membership mismatch", m)) + + # 14. Unsupported entry with a non-U id. + m = copy.deepcopy(base) + first_unsupported(m)["id"] = "A99" + mutations.append(("unsupported entry with non-U id", m)) + + return mutations + + +def main() -> int: + if not VALIDATOR.is_file(): + print(f"FAIL: validator not found at {VALIDATOR}", file=sys.stderr) + return 1 + if not MANIFEST.is_file(): + print(f"FAIL: manifest not found at {MANIFEST}", file=sys.stderr) + return 1 + + failures: list[str] = [] + + # The real committed manifest must pass. + real_rc = run_validator(MANIFEST) + if real_rc != 0: + failures.append(f"real manifest should PASS (exit 0), got exit {real_rc}") + else: + print("ok: real manifest passes the validator") + + base = load_real_manifest() + mutations = build_mutations(base) + + with tempfile.TemporaryDirectory(prefix="tc-cuda-subset-selftest-") as tmp: + tmp_root = pathlib.Path(tmp) + for index, (label, mutated) in enumerate(mutations): + path = tmp_root / f"mutation_{index:02d}.json" + path.write_text(json.dumps(mutated, indent=2), encoding="utf-8") + rc = run_validator(path) + if rc == 0: + failures.append(f"mutation {label!r} should be REJECTED (exit 1), got exit 0") + else: + print(f"ok: rejected {label!r}") + + if failures: + print("tc-cuda subset validator selftest FAILED:", file=sys.stderr) + for item in failures: + print(f" - {item}", file=sys.stderr) + return 1 + + print(f"tc-cuda subset validator selftest OK: real manifest passed, " + f"{len(mutations)} mutations rejected") + return 0 + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/scripts/tc_cuda.py b/scripts/tc_cuda.py new file mode 100644 index 0000000..b2c7b89 --- /dev/null +++ b/scripts/tc_cuda.py @@ -0,0 +1,460 @@ +#!/usr/bin/env python3 +"""tc-cuda frontend CLI — checked, not lowered. + +This is the first real tc-cuda frontend slice. It is a Python-stdlib-only +tool that: + +1. Loads and validates ``docs/tc-cuda/subset.v1.json`` via the existing + ``scripts/check_tc_cuda_subset.py`` validator (single authority). +2. Lexes CUDA source, ignoring comments and string/char literals. +3. Rejects every unsupported construct with a named diagnostic containing + the manifest ``id`` and ``name``; exits non-zero. +4. Rejects unknown CUDA double-underscore intrinsics/qualifiers not + represented in the supported authority. +5. On accepted source, emits a deterministic kernel manifest listing the + discovered ``__global__`` kernels with status ``checked`` (never + ``lowered``). + +Usage: + python3 scripts/tc_cuda.py check SOURCE.cu [--manifest-output PATH] + +Exit codes: + 0 accepted (source is within the supported subset) + 1 rejected (one or more unsupported/unknown constructs found) + 2 usage or authority error +""" + +from __future__ import annotations + +import argparse +import json +import pathlib +import re +import sys +from typing import Any + +# --------------------------------------------------------------------------- +# Authority loading — reuse the existing validator, do not duplicate policy. +# --------------------------------------------------------------------------- + +ROOT = pathlib.Path(__file__).resolve().parents[1] +DEFAULT_MANIFEST = ROOT / "docs" / "tc-cuda" / "subset.v1.json" + + +def _load_authority(manifest_path: pathlib.Path) -> dict[str, Any]: + """Load and validate the subset authority via check_tc_cuda_subset.""" + import importlib.util + + validator_path = ROOT / "scripts" / "check_tc_cuda_subset.py" + spec = importlib.util.spec_from_file_location( + "check_tc_cuda_subset", validator_path + ) + if spec is None or spec.loader is None: + raise SystemExit( + f"cannot import validator at {validator_path}" + ) + mod = importlib.util.module_from_spec(spec) + spec.loader.exec_module(mod) + + data = mod.load_manifest(manifest_path) + if not isinstance(data, dict): + raise SystemExit("subset manifest top-level must be a JSON object") + + # Run the validator's full check; if it fails, exit 2. + errors: list[str] = [] + mod.check_top_level(errors, data) + counts = mod.check_category_counts(errors, data) + supported_ids = mod.check_supported(errors, data, counts) + unsupported_ids = mod.check_unsupported(errors, data) + mod.check_duplicates(errors, supported_ids, unsupported_ids) + if errors: + print("tc-cuda: subset authority validation failed:", file=sys.stderr) + for e in errors: + print(f" - {e}", file=sys.stderr) + raise SystemExit(2) + + return data + + +# --------------------------------------------------------------------------- +# Lexer — strip comments and string/char literals, track line numbers. +# --------------------------------------------------------------------------- + + +def lex_cuda_source(text: str) -> str: + """Remove comments and string/char literals, preserving line structure. + + Returns a string where: + - ``//`` line comments are replaced with spaces (same line). + - ``/* ... */`` block comments are replaced with spaces (newlines kept). + - ``"..."`` string literals are replaced with a single space. + - ``'...'`` char literals are replaced with a single space. + - Backslash escapes inside strings/chars are handled. + + The resulting string has the same number of lines as the input, so + ``line = result.count('\\n', 0, pos) + 1`` gives the correct line number. + """ + out: list[str] = [] + i = 0 + n = len(text) + while i < n: + c = text[i] + # Line comment + if c == "/" and i + 1 < n and text[i + 1] == "/": + while i < n and text[i] != "\n": + out.append(" ") + i += 1 + # Block comment + elif c == "/" and i + 1 < n and text[i + 1] == "*": + out.append(" ") + out.append(" ") + i += 2 + while i < n: + if text[i] == "*" and i + 1 < n and text[i + 1] == "/": + out.append(" ") + out.append(" ") + i += 2 + break + out.append(text[i] if text[i] == "\n" else " ") + i += 1 + # String literal + elif c == '"': + out.append(" ") + i += 1 + while i < n and text[i] != '"': + if text[i] == "\\" and i + 1 < n: + i += 2 + else: + i += 1 + if i < n: + i += 1 # closing quote + # Char literal + elif c == "'": + out.append(" ") + i += 1 + while i < n and text[i] != "'": + if text[i] == "\\" and i + 1 < n: + i += 2 + else: + i += 1 + if i < n: + i += 1 # closing quote + else: + out.append(c) + i += 1 + return "".join(out) + + +def line_at(text: str, pos: int) -> int: + return text.count("\n", 0, pos) + 1 + + +# --------------------------------------------------------------------------- +# Unsupported-construct token extraction. +# +# Each unsupported manifest entry has a human-readable `name`. We derive +# concrete source-level tokens from that name. The token list is the +# detection surface; the diagnostic always uses the manifest `id` and `name`. +# --------------------------------------------------------------------------- + +# Map of unsupported id -> list of source tokens that indicate its presence. +# Tokens are matched as whole identifiers or as distinctive substrings. +_UNSUPPORTED_TOKENS: dict[str, list[str]] = { + "U1": ["wmma", "mma_sync", "mma."], + "U2": ["asm volatile", "asm("], + "U3": ["cp.async", "cp_async"], + "U4": ["ldmatrix"], + "U5": ["cooperative_groups", "cg::", "cooperative::"], + "U6": ["cudaGraph", "cudaGraphExec"], + "U7": ["cudaLaunchDevice", "cudaDeviceEnablePeerAccess"], + "U8": ["texture", "surface", "tex1D", "tex2D", "tex3D", "texFetch"], + "U9": ["__constant__"], + "U10": ["__ballot_sync", "__any_sync", "__all_sync"], + "U11": ["__syncwarp"], + "U12": ["__ldg"], + "U13": ["atomicAdd" ], # shared-memory atomics: atomicAdd on __shared__ + "U14": ["atomicAdd", "atomicSub", "atomicAnd", "atomicOr", "atomicXor"], + "U15": ["atomicCAS", "atomicExch", "atomicMax", "atomicMin"], + "U16": ["double", "fma(", "sqrt(", "exp(", "log(", "tanh(", "pow(", "sin(", "cos(", "floor(", "fmax(", "fmin(", "fabs("], + "U17": ["printf"], + "U18": ["malloc", "free", "assert"], + "U19": [], # recursion is structural; not detectable by token scan + "U20": ["virtual", "dynamic_cast", "typeid"], + "U21": ["template None: + self.code = code + self.message = message + self.line = line + + def render(self, source_path: str) -> str: + return f"tc-cuda: {self.code} {self.message} at {source_path}:{self.line}" + + +# --------------------------------------------------------------------------- +# Source scanning +# --------------------------------------------------------------------------- + + +def _find_token_positions(lexed: str, token: str) -> list[int]: + """Find all positions of `token` in `lexed` as a whole identifier or + as a distinctive substring. Returns a list of character positions.""" + positions: list[int] = [] + if " " in token: + # Multi-word token like "asm volatile" — simple substring search. + start = 0 + while True: + idx = lexed.find(token, start) + if idx == -1: + break + positions.append(idx) + start = idx + 1 + return positions + # Single-token: match as whole identifier (bounded by non-identifier chars). + pattern = re.compile( + r"(? list[Diagnostic]: + """Scan lexed source for unsupported constructs; return diagnostics.""" + diags: list[Diagnostic] = [] + for entry in unsupported_entries: + entry_id = entry["id"] + name = entry["name"] + tokens = _UNSUPPORTED_TOKENS.get(entry_id, []) + for token in tokens: + for pos in _find_token_positions(lexed, token): + line = line_at(lexed, pos) + diags.append( + Diagnostic( + code=entry_id, + message=f"unsupported construct {name!r} (token {token!r})", + line=line, + ) + ) + break # one diagnostic per entry per token is enough + return diags + + +def scan_unknown_intrinsics( + lexed: str, + source_path: str, +) -> list[Diagnostic]: + """Reject any __-prefixed identifier not in the supported authority.""" + diags: list[Diagnostic] = [] + # Match the full double-underscore-prefixed identifier greedily so that + # names such as ``__my_unknown__`` are reported in full (not truncated). + pattern = re.compile(r"__[A-Za-z][A-Za-z0-9_]*") + seen: set[str] = set() + for m in pattern.finditer(lexed): + ident = m.group(0) + # Normalise: strip trailing __ if present (e.g. __global__ -> __global__) + if ident not in seen: + seen.add(ident) + if ident in _SUPPORTED_DOUBLE_UNDERSCORE: + continue + if ident in _KNOWN_KEYWORDS: + continue + line = line_at(lexed, m.start()) + diags.append( + Diagnostic( + code="UNKNOWN", + message=f"unknown CUDA intrinsic/qualifier {ident!r} is not in the supported subset authority", + line=line, + ) + ) + return diags + + +def discover_global_kernels(lexed: str) -> list[str]: + """Find __global__ kernel function names in lexed source.""" + kernels: list[str] = [] + # Match: __global__ [qualifiers] return_type name ( + pattern = re.compile( + r"__global__\s+" + r"(?:static\s+)?(?:inline\s+)?(?:const\s+)?" + r"(?:void|[A-Za-z_][A-Za-z0-9_]*)\s+" + r"([A-Za-z_][A-Za-z0-9_]*)\s*\(" + ) + for m in pattern.finditer(lexed): + kernels.append(m.group(1)) + return sorted(set(kernels)) + + +# --------------------------------------------------------------------------- +# Manifest emission +# --------------------------------------------------------------------------- + + +def build_manifest( + source_path: str, + authority: dict[str, Any], + kernels: list[str], +) -> dict[str, Any]: + """Build a deterministic kernel manifest. Status is always 'checked'. + + The manifest is source-location independent: it records only the + authority identity and the discovered kernel names, never the absolute + path of the input file, so repeated runs over the same source produce + byte-identical output. + """ + return { + "schema": "tensorcore.tc-cuda.kernel-manifest.v1", + "authority": { + "schema": authority.get("schema"), + "version": authority.get("version"), + "total": authority.get("total"), + }, + "kernels": [ + {"name": k, "status": "checked"} for k in kernels + ], + "status": "checked", + "note": ( + "checked means the source was parsed and validated against the " + "subset authority. It does NOT mean the kernel was translated to " + "a backend or can execute. tc-cuda does not claim any backend " + "code generation has occurred." + ), + } + + +def emit_manifest(manifest: dict[str, Any], output_path: pathlib.Path) -> None: + text = json.dumps(manifest, indent=2, sort_keys=True) + "\n" + output_path.write_text(text, encoding="utf-8") + + +# --------------------------------------------------------------------------- +# CLI +# --------------------------------------------------------------------------- + + +def cmd_check(args: argparse.Namespace) -> int: + source_path = pathlib.Path(args.source) + if not source_path.is_file(): + print(f"tc-cuda: source file not found: {source_path}", file=sys.stderr) + return 2 + + authority = _load_authority(args.manifest) + + raw = source_path.read_text(encoding="utf-8") + lexed = lex_cuda_source(raw) + + unsupported_entries = authority.get("unsupported", []) + diags = scan_unsupported(lexed, str(source_path), unsupported_entries) + diags.extend(scan_unknown_intrinsics(lexed, str(source_path))) + + if diags: + for d in sorted(diags, key=lambda d: (d.line, d.code)): + print(d.render(str(source_path)), file=sys.stderr) + print( + f"tc-cuda: {len(diags)} diagnostic(s); source rejected.", + file=sys.stderr, + ) + return 1 + + kernels = discover_global_kernels(lexed) + manifest = build_manifest(source_path, authority, kernels) + + if args.manifest_output: + out = pathlib.Path(args.manifest_output) + emit_manifest(manifest, out) + print(f"tc-cuda: accepted; manifest written to {out}") + print(f"tc-cuda: {len(kernels)} __global__ kernel(s) discovered: {kernels}") + else: + print(f"tc-cuda: accepted; {len(kernels)} __global__ kernel(s): {kernels}") + + return 0 + + +def main() -> int: + parser = argparse.ArgumentParser( + description="tc-cuda frontend — checked, not lowered." + ) + sub = parser.add_subparsers(dest="command", required=True) + + p_check = sub.add_parser( + "check", + help="Check a CUDA source file against the tc-cuda v1 subset authority.", + ) + p_check.add_argument("source", type=str, help="Path to the .cu source file.") + p_check.add_argument( + "--manifest-output", + type=str, + default=None, + help="If set, write the kernel manifest JSON to this path.", + ) + p_check.add_argument( + "--manifest", + type=pathlib.Path, + default=DEFAULT_MANIFEST, + help="Path to subset.v1.json (default: docs/tc-cuda/subset.v1.json).", + ) + p_check.set_defaults(func=cmd_check) + + args = parser.parse_args() + return args.func(args) + + +if __name__ == "__main__": + raise SystemExit(main()) diff --git a/scripts/tc_cuda_selftest.py b/scripts/tc_cuda_selftest.py new file mode 100644 index 0000000..085df10 --- /dev/null +++ b/scripts/tc_cuda_selftest.py @@ -0,0 +1,300 @@ +#!/usr/bin/env python3 +"""Selftest for the tc-cuda frontend CLI. + +Generates fixtures covering: +- every one of the 25 unsupported manifest entries (rejected with id+name), +- comment and string-literal false positives (accepted), +- unknown __-prefixed intrinsic rejection, +- accepted minimal and rmsnorm-shaped kernels, +- manifest determinism (byte-for-byte identical across two runs). + +Usage: + python3 scripts/tc_cuda_selftest.py + +Exit 0 on success, 1 on any failure. +""" + +from __future__ import annotations + +import json +import pathlib +import subprocess +import sys +import tempfile +import textwrap + +ROOT = pathlib.Path(__file__).resolve().parents[1] +TC_CUDA = ROOT / "scripts" / "tc_cuda.py" +SUBSET = ROOT / "docs" / "tc-cuda" / "subset.v1.json" + +PASS = 0 +FAIL = 0 + + +def report(ok: bool, label: str, detail: str = "") -> None: + global PASS, FAIL + if ok: + PASS += 1 + print(f" PASS {label}") + else: + FAIL += 1 + print(f" FAIL {label}") + if detail: + for line in detail.splitlines(): + print(f" {line}") + + +def run_tc_cuda(source: str, manifest_output: str | None = None) -> tuple[int, str, str]: + """Write source to a temp .cu file, run tc_cuda.py check, return (rc, stdout, stderr).""" + with tempfile.NamedTemporaryFile( + mode="w", suffix=".cu", delete=False, encoding="utf-8" + ) as f: + f.write(source) + path = f.name + + cmd = [sys.executable, str(TC_CUDA), "check", path] + if manifest_output: + cmd += ["--manifest-output", manifest_output] + + proc = subprocess.run(cmd, capture_output=True, text=True, timeout=30) + pathlib.Path(path).unlink(missing_ok=True) + return proc.returncode, proc.stdout, proc.stderr + + +def load_unsupported_entries() -> list[dict]: + data = json.loads(SUBSET.read_text(encoding="utf-8")) + return data.get("unsupported", []) + + +def test_unsupported_entries() -> None: + """Each of the 25 unsupported entries must be rejected with its id and name.""" + print("\n[unsupported entries — 25 fixtures]") + entries = load_unsupported_entries() + assert len(entries) == 25, f"expected 25 unsupported entries, got {len(entries)}" + + # For each entry, generate a minimal CUDA source that contains a token + # that the scanner will detect. We use the first token from the + # _UNSUPPORTED_TOKENS map for that id. + import importlib.util + spec = importlib.util.spec_from_file_location("tc_cuda", TC_CUDA) + mod = importlib.util.module_from_spec(spec) + spec.loader.exec_module(mod) + + for entry in entries: + eid = entry["id"] + name = entry["name"] + tokens = mod._UNSUPPORTED_TOKENS.get(eid, []) + if not tokens: + # U19 (recursion) is structural; skip token-based test. + report(True, f"{eid} {name[:40]} (structural, skipped)") + continue + token = tokens[0] + # Build a source that uses this token in a __global__ kernel. + if " " in token: + # Multi-word token like "asm volatile" + src = textwrap.dedent(f"""\ + __global__ void k() {{ + {token} {{}}; + }} + """) + else: + # Single token: use it as a function call or identifier. + if token.startswith("__"): + src = textwrap.dedent(f"""\ + __global__ void k() {{ + {token}(); + }} + """) + elif token.startswith("cuda"): + src = textwrap.dedent(f"""\ + __global__ void k() {{ + {token}(0); + }} + """) + elif token.endswith("."): + # e.g. "mma." — use as namespace + src = textwrap.dedent(f"""\ + __global__ void k() {{ + {token}foo(); + }} + """) + else: + src = textwrap.dedent(f"""\ + __global__ void k() {{ + {token}(0); + }} + """) + + rc, out, err = run_tc_cuda(src) + # Must be rejected (non-zero exit) and diagnostic must contain id and name. + ok = rc != 0 and eid in err + report(ok, f"{eid} {name[:50]}", f"rc={rc} stderr={err[:200]}") + + +def test_comment_string_false_positives() -> None: + """Unsupported tokens inside comments or strings must NOT trigger rejection.""" + print("\n[comment/string false positives]") + src = textwrap.dedent("""\ + // __ballot_sync is not supported + /* __syncwarp __ldg __launch_bounds__ */ + __global__ void k() { + const char* s = "__ballot_sync __syncwarp"; + const char c = 'x'; + } + """) + rc, out, err = run_tc_cuda(src) + report(rc == 0, "comment/string false positives accepted", f"rc={rc} stderr={err[:200]}") + + +def test_unknown_intrinsic() -> None: + """Unknown __-prefixed intrinsic must be rejected.""" + print("\n[unknown intrinsic rejection]") + src = textwrap.dedent("""\ + __global__ void k() { + __my_unknown_intrinsic(); + } + """) + rc, out, err = run_tc_cuda(src) + ok = rc != 0 and "__my_unknown_intrinsic" in err + report(ok, "unknown __ intrinsic rejected", f"rc={rc} stderr={err[:200]}") + + +def test_accepted_minimal() -> None: + """A minimal accepted kernel with only supported constructs.""" + print("\n[accepted minimal kernel]") + src = textwrap.dedent("""\ + __device__ float helper(float x) { + return x * 2.0f; + } + + __global__ void add(float* a, float* b, float* out, int n) { + int i = threadIdx.x + blockIdx.x * blockDim.x; + if (i < n) { + out[i] = helper(a[i]) + b[i]; + } + } + """) + rc, out, err = run_tc_cuda(src) + report(rc == 0, "minimal kernel accepted", f"rc={rc} stderr={err[:200]}") + + +def test_accepted_rmsnorm() -> None: + """An rmsnorm-shaped kernel using only supported constructs.""" + print("\n[accepted rmsnorm-shaped kernel]") + src = textwrap.dedent("""\ + #include + + __device__ __half2float(__half h) { + return 0.0f; + } + + __global__ void rmsnorm_forward_kernel( + const __half* __restrict__ input, + __half* __restrict__ output, + const __half* __restrict__ weight, + int rows, + int cols, + float eps + ) { + int row = blockIdx.x; + if (row >= rows) return; + + const __half* in_row = input + (size_t)row * cols; + __half* out_row = output + (size_t)row * cols; + + float sum_sq = 0.0f; + for (int i = threadIdx.x; i < cols; i += blockDim.x) { + float v = __half2float(in_row[i]); + sum_sq += v * v; + } + + __shared__ float s_sum; + if (threadIdx.x == 0) s_sum = sum_sq; + __syncthreads(); + + float rms = rsqrtf(s_sum / (float)cols + eps); + + for (int i = threadIdx.x; i < cols; i += blockDim.x) { + float v = __half2float(in_row[i]); + out_row[i] = __float2half_rn(v * rms * __half2float(weight[i])); + } + } + """) + rc, out, err = run_tc_cuda(src) + report(rc == 0, "rmsnorm-shaped kernel accepted", f"rc={rc} stderr={err[:300]}") + + +def test_manifest_determinism() -> None: + """Two runs on the same source must produce byte-identical manifests.""" + print("\n[manifest determinism]") + src = textwrap.dedent("""\ + __global__ void k1(float* a) { a[0] = 1.0f; } + __global__ void k2(float* b) { b[0] = 2.0f; } + """) + + with tempfile.TemporaryDirectory() as td: + m1 = str(pathlib.Path(td) / "m1.json") + m2 = str(pathlib.Path(td) / "m2.json") + + rc1, out1, err1 = run_tc_cuda(src, manifest_output=m1) + rc2, out2, err2 = run_tc_cuda(src, manifest_output=m2) + + if rc1 != 0 or rc2 != 0: + report(False, "manifest determinism (both runs accepted)", + f"rc1={rc1} rc2={rc2} err1={err1[:200]} err2={err2[:200]}") + return + + b1 = pathlib.Path(m1).read_bytes() + b2 = pathlib.Path(m2).read_bytes() + ok = b1 == b2 + report(ok, "manifest byte-identical across two runs", + "" if ok else f"m1={b1[:200]} m2={b2[:200]}") + + # Also verify the manifest content is correct. + data = json.loads(b1) + kernels = [k["name"] for k in data.get("kernels", [])] + ok2 = sorted(kernels) == ["k1", "k2"] and data.get("status") == "checked" + report(ok2, "manifest content correct (kernels + status=checked)", + f"kernels={kernels} status={data.get('status')}") + + +def test_manifest_status_never_lowered() -> None: + """The manifest must never contain status 'lowered'.""" + print("\n[manifest status never lowered]") + src = "__global__ void k() { }" + with tempfile.TemporaryDirectory() as td: + m = str(pathlib.Path(td) / "m.json") + rc, out, err = run_tc_cuda(src, manifest_output=m) + if rc != 0: + report(False, "accepted for manifest check", f"rc={rc} err={err[:200]}") + return + text = pathlib.Path(m).read_text(encoding="utf-8") + ok = "lowered" not in text + report(ok, "manifest does not contain 'lowered'", + "" if ok else text[:300]) + + +def main() -> int: + print("tc-cuda frontend selftest") + print(f" tc_cuda.py: {TC_CUDA}") + print(f" subset: {SUBSET}") + + test_unsupported_entries() + test_comment_string_false_positives() + test_unknown_intrinsic() + test_accepted_minimal() + test_accepted_rmsnorm() + test_manifest_determinism() + test_manifest_status_never_lowered() + + print(f"\n{'=' * 50}") + print(f" PASS: {PASS} FAIL: {FAIL}") + if FAIL > 0: + print(" RESULT: FAIL") + return 1 + print(" RESULT: PASS") + return 0 + + +if __name__ == "__main__": + raise SystemExit(main())