## 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>
## 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>
## 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#540Closes#541🤖 Generated with [Claude Code](https://claude.com/claude-code)
---------
Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
Ang Gao