forked from ccf-ai-infra/TileOPs-Metax
4 Commits
| Author | SHA1 | Message | Date |
|---|---|---|---|
wawahejun
|
ce1b15c28f |
docs: add MetaX C500 summer-camp guide on top of upstream dev
Source tree is identical to MetaX-MACA/TileOPs-Metax dev at f02d3d8; this commit carries only the summer-camp documentation and PR templates. Content verified by running everything on a real MetaX C500 (MACA 3.7.1.5, torch 2.8.0+metax3.7.1.3, tilelang 0.1.10+cuda.gitf549117c, sGPU slice 16000 MiB). Installation (high severity). The documented `make install`, `pip install tileops`, and bare `python3 -m venv .venv` steps destroy a working MACA environment. The container's TileLang is an in-place source build imported via PYTHONPATH, so pip reports it as absent and resolves the official CUDA wheel over it; a venv without --system-site-packages cuts off the MetaX PyTorch build and the ABI-coupled apache-tvm-ffi. Replace those steps with the PYTHONPATH setup, document that tileops needs no install at all, and note that --no-deps is the only safe install form (as scripts/ci/install_tileops.sh already does). Flag -c constraints.txt as CUDA-CI-only for the same ABI reason. Add TileLang provenance and backend checks to the verification list, which previously covered mx-smi, torch and einops but not the component most likely to be wrong. Quick start. GemmOp(M, N, K, dtype=...) does not match the implementation -- GemmOp is input-inferred and takes only trans_a/trans_b. Fix the signature and document the trans_b default, keeping the original M,N,K of 1024,1024,512, which passes on C500 via the MACA dispatch path. New sections 1.2 and 1.3 in the migration guide. Document the is_maca() dispatch to *_maca.py kernels, and that a gated kernel does not imply an unusable Op: GemmKernel declares [89, 90] and is gated on C500, yet GemmOp works because it dispatches to gemm_maca.py ([80, 86, 89, 90]). Availability must be judged from what the Op layer dispatches to, not from one kernel's supported_archs. List the 20 declarations that exclude 80 as unsuitable migration targets, and note that adding a *_maca.py kernel plus dispatch is a good target instead. Record that get_sm_version() reuses NVIDIA's encoding, so C500 reports 80 while sharing nothing with Ampere, and that the raw "architecture 80" message names no device. Document that a usable Op still has shape limits: SoftmaxFwdOp fails above a 1024-wide reduction dimension (mcErrorInvalidValue), independent of row count. Document that a parent process which has imported tilelang will see any subprocess that imports it again SIGKILLed with no output, which aborts tests/test_validate_manifest.py at exit 137, and give the deselect workaround. Roofline. Record the sGPU slice quota and state whether peaks are whole-card or slice-scaled; dividing a slice measurement by a whole-card peak yields an unexplainable efficiency. Verified on C500 against this tree: validate_manifest.py exit 0; 29 passed across test_ops_manifest.py, test_kernel_map_install.py and benchmarks/tests; GemmOp passes at 1024x1024x512, 1024^3 and 4096^3; the documented quick-start snippet and every self-check command run as written. pre-commit and ruff are unavailable in this container (installing them would invoke pip dependency resolution), so formatting was checked via git diff --check and end-of-file newlines instead. Squashed documentation commits by Beckylu <648245013@qq.com> and FrRay <1077376663@qq.com> covering the summer-camp guide, PR templates, and README translations. Co-Authored-By: Beckylu <648245013@qq.com> Co-Authored-By: FrRay <1077376663@qq.com> |
|
|
Cao Ying
|
1d99d4cda2
|
[CI] retire run-local venv across CI; install via image-baked stack (#1606)
## What
Full RFC §4 cleanup on top of the ephemeral containerized-runner
migration. The runner image bakes tilelang + the runtime/dev stack, so
CI stops building/copying per-run venvs and installs only tileops
(`--no-deps`) via `scripts/ci/install_tileops.sh`.
- **gpu-smoke**: drop `Set up Python`, `Resolve runtime state` (venv
hash/copy/mtime-sync/divergence guard), `Cleanup isolated fork state`;
install via `install_tileops.sh`; run with the image's `python3`;
trust-route PRs by collaborator permission (write/maintain/admin →
resident `nightly` pool; everyone else + lookup failure → `fork` pool,
fail-closed); reclaim the `/ci-cache` layout (no wheels, no tool-cache
prune); per-test `--timeout` + a `timeout-minutes` backstop so a wedged
kernel cannot hold the single runner.
- **nightly**: `setup_nightly_venv.sh` → `install_tileops.sh`; system
`python3`; cache env `/data7` → `/ci-cache`; per-test timeout on the
pytest runs.
- **runner-maintenance**: drop the retired `venv` runs-on label (→
`nightly`).
- **preflight**: pin CPU installs with `constraints.txt`.
- **pyproject / constraints**: tilelang as a compatibility range
(constraints pins its deps, not tilelang itself); pin `pytest-timeout`.
- **reclaim action + verify script**: `/ci-cache` layout, drop
`WHEEL_DIR`; trim stamp on the persistent cache.
- **delete**: `setup_nightly_venv.sh`, `ci_venv_hash.py` and their
obsolete tests.
## In scope vs follow-up
To keep CI dispatching, installing and importing cleanly on the new
stack, this PR DOES adapt the import surface: `gqa_fwd_fp8` guards its
`from tvm import tir` with a sentinel that raises a targeted error only
when the kernel is built, and the fp8-GQA / topk-selector
kernel-building smoke cases are skipped via a `tvm.tir` availability
gate (with a focused test for the gate). So **CI is green at import**,
not red.
**Out of scope (follow-up):** the actual kernel migration off `tvm.tir`
(`tir.call_extern` → `tilelang.language` `T.*`) and any other new-stack
kernel regressions surfaced by gpu-smoke (e.g. a build hang in
`FP8LightingIndexerKernel`). Those are tracked separately; until they
land, the corresponding smoke cases are skipped or fail on the per-test
timeout rather than wedging the runner.
---------
Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com>
|
|
|
Cao Ying
|
fefcb10a8b
|
[CI] Reproducible multi-stage runner image + single dependency source (#1600)
Closes #1599 ## Summary - Repo-side foundation of the CI dependency / cache redesign — additive and files-only; nothing consumes these files yet (the workflow cutover is a separate follow-up). - `constraints.txt`: single source of exact version pins for the CI/runner stack (torch, triton, apache-tvm-ffi, and tilelang runtime deps). - `scripts/ci/install_tileops.sh`: requires `tilelang` already present, then `pip install -e . --no-deps -c constraints.txt`; fails clearly if tilelang is missing (so pip cannot drift torch / apache-tvm-ffi). - `.github/runner/Dockerfile`: rewritten as a multi-stage build from the public `nvidia/cuda:12.9.1-devel-ubuntu22.04` base (`runtime` builds python3.12 via the deadsnakes PPA → `post-fa3` → `fullstack` → `final`); bakes no TileOPs source and no runner credentials. ## Test plan - [x] AC-1: `constraints.txt`, `scripts/ci/install_tileops.sh`, and the rewritten multi-stage `.github/runner/Dockerfile` exist with the requested structure (stages runtime/post-fa3/fullstack/final, public CUDA 12.9.1 base, constraints-only Docker COPY, tilelang preflight, `--no-deps` install). - [x] AC-2: `scripts/ci/install_tileops.sh` passes `shellcheck` (0.11.0, no diagnostics). - [x] AC-3: `.github/runner/Dockerfile` passes `hadolint --failure-threshold error` (2.12.0; only warning/info findings). - [ ] AC-4: **DEFERRED-TO-MANUAL on a GPU build host.** The image is intentionally NOT built in CI by this issue. Manual validation (image builds; `torch.version.cuda == "12.9"`; `tilelang` imports; cuBLAS matmul/bmm/einsum probe passes; `pytest -m smoke` passes) is performed by a maintainer on a host with an NVIDIA GPU + nvcc. Reviewers should not expect a CI image build here. --------- Co-authored-by: Ibuki 🍃 — a wind born from GPTs <Ibuki-wind@users.noreply.github.com> |
|
|
ChongLi
|
4da2dbba1d
|
[CI] Run nightly inside containerized runner (#1017)
Closes #1016 ## Summary - Move nightly workflow execution to assume it already runs inside the containerized self-hosted runner. - Remove workflow-level Docker startup, workspace ownership fixes, and the local tileops-runner:latest wrapper. - Add a nightly runner verification script for cache env vars, cache writability, and GPU visibility. ## Test plan - [x] pre-commit passed during git commit - [x] bash -n scripts/ci/verify_nightly_runner.sh - [x] python -c "import yaml; yaml.safe_load(open('.github/workflows/nightly.yml')); print('yaml ok')" - [x] git diff --check -- .github/workflows/nightly.yml scripts/ci/verify_nightly_runner.sh - [x] rg -n "docker run|tileops-runner:latest|Fix workspace ownership|/data/ci-cache" .github/workflows/nightly.yml returns no matches - [ ] actionlint .github/workflows/nightly.yml (not run locally; actionlint is not installed) ## Additional context This PR intentionally keeps the existing runner labels unchanged and does not modify runner-maintenance.yml. Cache cleanup policy is left for a separate follow-up, as scoped in #1016. |