forked from ccf-ai-infra/TileOPs-Metax
4 Commits
| Author | SHA1 | Message | Date |
|---|---|---|---|
Ang Gao
|
89708ee424
|
[Fix][Benchmark] Add serial validation to correct noisy autotune configs (#773)
## Summary
- Parallel warmup (`-n 8`) causes autotuner to select suboptimal configs
due to GPU contention noise, which get cached and reused by serial
benchmarks
- Add a Phase 2 serial validation pass in `warmup_kernel_cache.py` that
re-tunes on a quiet GPU with `.so` cache hits (compilation instant, only
profiling re-runs)
- Add autotune config column to `BenchmarkReport` so nightly reports
show which config was selected
## How it works
| Phase | Compilation cache (`.so`) | Autotune cache (`best_config`) |
|---|---|---|
| Phase 1: parallel warmup (`-n 8`) | Write (parallel, fast) | Write
(noisy — GPU contention) |
| Phase 2: serial validation | Read (cache hit, instant) | Overwrite
(accurate — exclusive GPU) |
`conftest_warmup.py` supports a new `TILEOPS_WARMUP_VALIDATE=1` mode
that patches `AutoTuner._load_result_from_disk` to return `None`,
forcing re-tune while preserving `.so` cache.
## Evidence
Same kernel (`avg_pool2d vision-3x3-s2`) selects different configs under
contention vs serial:
- **Parallel (8 workers)**: `{'block_m': 128, 'block_c': 64, 'threads':
256}`
- **Serial (3/3 runs consistent)**: `{'block_m': 64, 'block_c': 32,
'threads': 256}`
Fixes #772
## Test plan
- [x] Full warmup + validation pass completes on all benchmark files
- [x] Serial validation correctly overwrites noisy autotune cache
- [x] `BenchmarkReport` includes config column for tileops entries
- [x] Verified on H200 with `tileops-runner:latest`
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
|
|
|
Ang Gao
|
8bccfa30e7
|
[CI] Fix warmup OOM, enable autotuner caching, skip baselines (#685)
## Summary Three fixes to the nightly warmup phase: 1. **Fix GPU OOM causing warmup to hang** — pytest-xdist workers accumulated GPU memory over their lifetime; idle workers held ~140GB while the last worker starved and hung at 98%. Added `gc.collect()` + `torch.cuda.empty_cache()` after each test. 2. **Enable autotuner result caching** — Removed `TILELANG_AUTO_TUNING_DISABLE_CACHE=1` and the `do_bench` mock patch. Warmup now runs real GPU profiling and saves `best_config.json` to disk. Subsequent warmup and benchmark runs hit the autotuner cache (~2ms per kernel instead of minutes). 3. **Skip baseline profiling during warmup** — Patched `BenchmarkBase.profile` to return dummy results for non-Op functors (torch-cublas, FA3, etc.). Baseline measurements serve no purpose during warmup. Also bumped default pytest-xdist workers from 8 to 16. **Result**: warmup dropped from 38min (hanging at 98%) to 9min (912/912 passed). ## Changes - `scripts/warmup_kernel_cache.py` — Remove `TILELANG_AUTO_TUNING_DISABLE_CACHE=1`, default `-n 16` - `scripts/conftest_warmup.py` — Remove `do_bench` mock, add GPU memory cleanup hook, add baseline skip - `.github/workflows/nightly.yml` — Bump warmup to `-n 16` ## Test plan - [x] Warmup completes without OOM (912/912 passed, GPU memory stable at ~20-30GB) - [x] Warmup time: 38min (hanging) → 9min with caches warm - [x] Autotuner cache populated (best_config.json written to disk) - [ ] Verify all 4 nightly phases complete via `workflow_dispatch` 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com> |
|
|
Ang Gao
|
27c6688808
|
[BugFix][CI] Fix warmup error handling, monkeypatch cleanup, and portable cache paths (#548)
## Summary
- Fix warmup script to propagate infrastructure errors (pytest exit
codes 2-5) while still tolerating test failures (exit code 1)
- Restore original `ThreadPoolExecutor` in `pytest_unconfigure` to
prevent global state leaks
- Replace all hardcoded `/home/ci-runner/` paths with `${HOME}/` for
runner portability
Closes #545
Closes #546
## Changes
### 1. Warmup exit code handling (#545)
```diff
- sys.exit(0) # unconditional
+ if exit_code in (0, 1):
+ sys.exit(0) # tests passed or failed — compilation succeeded
+ else:
+ sys.exit(exit_code) # infrastructure error — surface it
```
| Exit Code | Meaning | Before | After |
|-----------|---------|--------|-------|
| 0 | All passed | exit 0 | exit 0 |
| 1 | Test failures | exit 0 | exit 0 (compilation still worked) |
| 2 | Interrupted / bad args | exit 0 (swallowed) | **exit 2** |
| 3 | Internal error | exit 0 (swallowed) | **exit 3** |
| 4 | Usage error | exit 0 (swallowed) | **exit 4** |
| 5 | No tests collected | exit 0 (swallowed) | **exit 5** |
### 2. ThreadPoolExecutor cleanup (#545)
```diff
def pytest_unconfigure(config):
patcher = getattr(config, "_warmup_patcher", None)
if patcher is not None:
patcher.stop()
+
+ orig_pool = getattr(config, "_warmup_orig_pool", None)
+ if orig_pool is not None:
+ concurrent.futures.ThreadPoolExecutor = orig_pool
```
### 3. Portable cache paths (#546)
```diff
- TILELANG_CACHE_DIR="/home/ci-runner/.tilelang/cache"
+ TILELANG_CACHE_DIR="${HOME}/.tilelang/cache"
```
Applied to all 4 jobs (cache-warmup, benchmark, op_test, packaging). 14
occurrences replaced.
## Test plan
- [ ] Verify warmup exits non-zero when pytest encounters infrastructure
errors
- [ ] Verify warmup exits 0 when tests fail but compilation completed
- [ ] Verify no hardcoded `/home/ci-runner/` paths remain
- [ ] Verify cache paths resolve correctly on CI runner
🤖 Generated with [Claude Code](https://claude.com/claude-code)
---------
Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
|
|
|
Ang Gao
|
7b94393188
|
[CI] Separate kernel compilation from profiling and parallelize compilation in nightly (#542)
## Summary - Add a `cache-warmup` job to the nightly workflow that pre-compiles all benchmark kernel variants before the benchmark job runs, separating compilation (CPU-bound) from profiling (GPU-bound) - Fix venv cache hash to only include dependency fields from `pyproject.toml`, preventing unnecessary rebuilds from linting/style config changes - Restructure nightly job flow: `cache-warmup → benchmark → op_test → packaging` ## Background In [nightly run #23161046010](https://github.com/tile-ai/TileOPs/actions/runs/23161046010), the benchmark job hit the 90-minute timeout at only 48% progress. Analysis showed ~80-90% of the time was spent on kernel compilation (`g++`/`cc1plus` still running at cancellation), not GPU profiling. Root causes: 1. **Venv prefix rename** in #530 caused `VENV_REUSED=false` → pip reinstalled tilelang → `libtilelang.so` mtime changed → **all kernel caches invalidated** (tilelang cache keys include `.so` mtime) 2. **`tune=True` benchmarks** (GEMM 288 configs, MHA, etc.) compile hundreds of kernel variants through TIR → CUDA → nvcc/g++ pipeline, all serially during pytest execution ## Changes ### 1. Kernel cache warmup (`scripts/warmup_kernel_cache.py`) - Monkeypatches `tilelang.profiler.do_bench` to return dummy value → compilation happens normally, GPU profiling skipped - Sets `TILELANG_AUTO_TUNING_DISABLE_CACHE=1` → dummy results not persisted to autotuner cache - Two levels of parallelism: - **pytest-xdist** (`-n 8`): 8 worker processes running different benchmark test cases simultaneously - **ThreadPoolExecutor** (`--max-workers 64`): within each `tune=True` op, autotuner compiles up to 64 config variants in parallel - **Zero maintenance**: new ops with benchmarks are automatically discovered and pre-compiled - Supports `--shard`/`--total-shards` for future parallelization across runners ### 2. Venv hash fix (Closes #540) ```diff - sha256sum pyproject.toml + python3 scripts/dep_hash.py # hashes only dependencies, optional-deps, requires-python, build-requires ``` Applied to all 4 jobs (cache-warmup, benchmark, op_test, packaging). Changes to ruff rules, codespell config, pytest markers etc. no longer trigger venv rebuild. ### 3. Job dependency restructure & timeout adjustments ``` cache-warmup (60 min) → benchmark (120 min) → op_test (180 min) → packaging ``` GPU jobs (benchmark, op_test) remain serial for measurement accuracy. Compilation cost is moved to the warmup phase. ## Local benchmark results (H200) Tested on a clean tilelang cache to simulate a worst-case cold start: | Phase | Duration | |-------|----------| | Warmup (n=8, --max-workers 64, cold cache) | 24 min | | Benchmark (warm cache, first run) | 73 min | | **Total** | **97 min** | ### Comparison with the previous approach | | Before | After (this PR) | |---|---|---| | Compilation | Serial, interleaved with profiling. 90 min timeout at 48% — estimated **3h+** to complete | Warmup job: 8 pytest workers × 64 compile threads, **24 min** | | Benchmark | Blocked behind compilation, never finished | All kernels pre-compiled, **73 min** (first run) | | **Total** | **> 3h (timeout, incomplete)** | **97 min (complete)** | ### Expected time in subsequent nightly runs (steady state) Both the kernel cache and autotuner result cache are stored in persistent directories on the self-hosted runner (`TILELANG_CACHE_DIR`). As long as the venv is not rebuilt (i.e., dependencies unchanged), these caches persist across nightly runs: | Phase | First run (cold cache) | Subsequent runs (warm cache) | |-------|----------------------|---------------------------| | Warmup | 24 min (full compilation) | **~2-3 min** (all cache hits, no actual compilation) | | Benchmark | 73 min (autotuner profiles all configs) | **~30-40 min** (autotuner cache hit → skip profiling 288 configs per op, only benchmark the best config) | | **Total** | **97 min** | **~35-45 min** | ### Additional notes - **Deepseek kernel issues**: Several Deepseek-related kernels currently have compilation or runtime issues (6 consecutive failures at ~48% progress). Once fixed, the time spent on their failed compilation/retry will be recovered. - **Warmup failures are non-blocking**: The warmup job uses `continue-on-error: true` and always exits 0. Any kernel that fails to compile during warmup will simply be compiled on-demand during the benchmark job (same behavior as before this PR). ## Test plan - [x] Verify warmup script compiles kernels and populates cache (local H200) - [x] Verify benchmark completes with warm cache (local H200, 73 min) - [x] Verify n=8 xdist parallelism works (24 min vs n=4 at 32 min) - [ ] Verify nightly completes within timeout on CI runner - [ ] Verify venv hash is stable across non-dependency pyproject.toml changes Closes #540 Closes #541 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com> |