forked from ccf-ai-infra/TileOPs-Metax
ce1b15c28f
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> |
||
|---|---|---|
| .. | ||
| moe | ||
| test_grouped_gemm_persistent.py | ||
| test_grouped_gemm_persistent_3wg.py | ||
| test_kernel_autotune_binding.py | ||
| test_moe_grouped_gemm_3wg_fused_act.py | ||