Skip to content
Draft
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
21 commits
Select commit Hold shift + click to select a range
9c60a30
refactor(infra): remove tir-bench agent skill (#2)
spectrometerHBH Jun 17, 2026
a42ad16
feat(lower-tirx): vector PTX ld/st and remove copy_xxb (#6)
spectrometerHBH Jun 24, 2026
cd0c540
fix(arith): add overflow-free uint32/uint64 floormod/floordiv rewrite…
spectrometerHBH Jun 24, 2026
6d10284
feat(infra): tir-bench helper, rounds, and microsecond timings (#8)
spectrometerHBH Jun 24, 2026
ef79330
feat(infra): TFLOAT32 descriptor support for copy_async(tma) (#9)
spectrometerHBH Jun 24, 2026
91e1506
feat(infra): plumb TVM_CUDA_PTXAS_EXTRA_OPTS through nvcc and nvrtc (…
spectrometerHBH Jun 24, 2026
3d46bf8
feat(op-dispatch): add dense fp8 and tf32 tcgen05 gemm_async paths (#11)
spectrometerHBH Jun 25, 2026
0235be7
fix(infra): reg-copy R2S/S2R for tcgen05 split-laneid atom layouts (#12)
spectrometerHBH Jun 26, 2026
9453a29
feat(op): add CUDA elementwise dispatch for maximum tile op (#13)
spectrometerHBH Jun 26, 2026
a96c6bf
fix(lower-tirx): elementwise reg slice under sctx.target (B00011) (#14)
spectrometerHBH Jun 26, 2026
84de3d9
refactor(op-dispatch): drop redundant pre-slice canonicalize in ld_st…
spectrometerHBH Jun 27, 2026
6f6af76
fix(op-dispatch): cache one tcgen05.cp descriptor, patch SMEM addr pe…
spectrometerHBH Jun 27, 2026
b310ba2
feat(op-dispatch): support uint32 shape extents in TMA copy (#17)
spectrometerHBH Jun 27, 2026
5cf4d40
perf(op-dispatch): direct ld/st.shared for shared->shared warp permut…
spectrometerHBH Jun 27, 2026
4a5b5ba
perf(op-dispatch): per-MMA SMEM descriptor with hoist/recompute in ge…
spectrometerHBH Jun 27, 2026
3f9ce07
fix(lower-tirx): support PrimType dtype handling (#21)
spectrometerHBH Jun 27, 2026
c732226
feat(infra): Triton-standard bench harness with proton default timer …
spectrometerHBH Jul 2, 2026
515d858
refactor(infra): always build and run all reference impls (#27)
spectrometerHBH Jul 4, 2026
cf11478
Remove dead async_structs, merge tile-primitive headers, rename Predi…
spectrometerHBH Jul 4, 2026
d6ef5bc
feat(layout): buffer dim-surgery views (unflatten/flatten/select/narr…
spectrometerHBH Jul 5, 2026
0f3c9d6
perf(kernel): optimize Blackwell grouped GEMM primitives
spectrometerHBH Jul 12, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
13 changes: 1 addition & 12 deletions .agents/scripts/monitor_gpu.sh
Original file line number Diff line number Diff line change
Expand Up @@ -23,18 +23,7 @@ while [[ $# -gt 0 ]]; do
--interval) INTERVAL="$2"; shift 2 ;;
--log) LOG="$2"; shift 2 ;;
-h|--help)
cat <<'EOF'
Watch a single GPU for foreign processes (anyone other than the current
user) appearing during a long-running test. Intended companion to
`/tir-test`: leave this running in a side terminal while pytest runs, and
it will alert if someone else lands on the same GPU.

Usage:
monitor_gpu.sh # uses $CUDA_VISIBLE_DEVICES, defaults to 0
monitor_gpu.sh --gpu 3 # watch GPU 3
monitor_gpu.sh --gpu 3 --interval 2 # poll every 2 seconds
monitor_gpu.sh --log /tmp/gpu.log # also tee to a log file
EOF
sed -n '2,12p' "$0" | sed 's/^# \{0,1\}//'

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

medium

Hardcoding line numbers like '2,12p' in sed is extremely fragile because any change to the top of the script (such as adding comments, changing imports, or adding license headers) will break the help output. Keeping the heredoc cat <<'EOF' is much more robust and standard.

Suggested change
sed -n '2,12p' "$0" | sed 's/^# \{0,1\}//'
cat <<'EOF'
Watch a single GPU for foreign processes (anyone other than the current
user) appearing during a long-running test. Intended companion to
'/tir-test': leave this running in a side terminal while pytest runs, and
it will alert if someone else lands on the same GPU.
Usage:
monitor_gpu.sh # uses $CUDA_VISIBLE_DEVICES, defaults to 0
monitor_gpu.sh --gpu 3 # watch GPU 3
monitor_gpu.sh --gpu 3 --interval 2 # poll every 2 seconds
monitor_gpu.sh --log /tmp/gpu.log # also tee to a log file
EOF

exit 0 ;;
*) echo "unknown arg: $1" >&2; exit 2 ;;
esac
Expand Down
195 changes: 0 additions & 195 deletions .agents/skills/tir-bench/SKILL.md

This file was deleted.

25 changes: 20 additions & 5 deletions .agents/skills/tir-test/SKILL.md
Original file line number Diff line number Diff line change
Expand Up @@ -2,9 +2,13 @@ Run the full TIRX test suite.

## Steps

1. Select the least busy GPU to avoid conflicts:
1. Install the kernel package and select the least busy GPU:
```bash
pip install -e /path/to/tirx-kernels-staging # or sibling tirx-kernels checkout
export CUDA_VISIBLE_DEVICES=$(nvidia-smi --query-gpu=index,memory.used --format=csv,noheader,nounits | sort -t',' -k2 -n | head -1 | cut -d',' -f1 | tr -d ' ')
export TVM_PATH=/path/to/tvm
export PYTHONPATH="${TVM_PATH}/python"
export TVM_LIBRARY_PATH="${TVM_PATH}/build/lib"
```

2. Start the GPU monitor in the background so we can detect if anyone else lands on the same GPU mid-run:
Expand All @@ -15,18 +19,29 @@ Run the full TIRX test suite.
trap 'kill $MON_PID 2>/dev/null' EXIT
```

3. Run the full test suite with xdist parallelism:
3. Import gate — bench workloads: fail fast if any kernel listed in `workloads.yaml` fails to import:
```bash
pytest tests/python/tirx/ -n auto
python -m tirx_kernels.bench_suite --check-imports
```
A non-zero exit means a pinned workload kernel failed to import — fix it before proceeding.

4. Stop the monitor and check for foreign GPU usage during the run:
4. Full kernel import gate (correctness test suite coverage):
```bash
python -m tirx_kernels.registry --cc 10 --strict
```

5. Run the full test suite with xdist parallelism:
```bash
pytest tests/python/tirx/ -n 16
```

6. Stop the monitor and check for foreign GPU usage during the run:
```bash
kill $MON_PID 2>/dev/null; wait $MON_PID 2>/dev/null
grep -E 'FOREIGN USER|\[FOREIGN\]' "$GPU_LOG" || echo "no foreign GPU usage observed"
```

5. Report results: total passed, failed, skipped, errors. If any foreign-user events are present in step 4, mention them — flaky failures should be re-evaluated on a clean GPU before being attributed to code changes.
7. Report results: total passed, failed, skipped, errors — and the import-gate results from steps 3–4. If any foreign-user events are present in step 6, mention them — flaky failures should be re-evaluated on a clean GPU before being attributed to code changes.

## Failure triage rules

Expand Down
5 changes: 0 additions & 5 deletions .gitignore
Original file line number Diff line number Diff line change
Expand Up @@ -183,7 +183,6 @@ cscope*
perf
.bash_history
*.json
!.claude/commands/tir-bench/baseline.json
*.params
*.ro
*.onnx
Expand Down Expand Up @@ -295,7 +294,3 @@ python/bin/
python/typing_extensions.py
python/*.dist-info/
pytest-of-bohanhou/

# tir-bench run artifacts (regenerable; see .claude/commands/tir-bench.md)
.tir-bench/
.tir-bench-*/
4 changes: 2 additions & 2 deletions AGENTS.md
Original file line number Diff line number Diff line change
Expand Up @@ -85,8 +85,8 @@ the branch instead of the whole repository:
pre-commit run --files <changed-file>...
```

Use `.agents/skills/tir-build`, `.agents/skills/tir-test`, and
`.agents/skills/tir-bench` when their workflows apply.
Use `python -m tirx_kernels.bench_suite` in the **tirx-kernels** repo
(`tirx_kernels/bench_suite/`) when that workflow applies.

## Coding Conventions

Expand Down
103 changes: 0 additions & 103 deletions include/tvm/tirx/async_structs.h

This file was deleted.

Loading