Repository navigation
[FEAT][platforms]: integrate persistent kernel Auto-Tuning into Hardware Abstraction Layer #50
Description
Description:
Context:
Optimal Triton/CUDA grid configurations (e.g., BLOCK_M, BLOCK_N, num_warps, num_stages) vary drastically between hardware generations (e.g., Ada Lovelace vs. Hopper vs. MI300X). Hardcoding these parameters leaves significant performance on the table.
Tasks:
- Implement an autotuning wrapper in the HAL (rl_engine/platforms/autotuner.py) that profiles operator performance on the first execution.
- Cache the best-performing configurations locally (e.g., to a .kernel_align_cache file) based on the specific GPU PCIe ID and shape of the input tensors.
- Refactor existing Triton operators like ops/triton/base.py to inherit from this autotuning wrapper rather than using static @triton.autotune heuristics.
Activity
fusheng-ji commentedon Oct 11, 2026
Hi, thanks for opening this issue! I wanted to let you know that I've worked on the CUDA side of it in #536, and it's ready for review whenever you have time.
In short, it adds an opt-in persistent autotuner in
rl_engine/platforms/autotuner.py, with a private on-disk cache keyed by GPU identity (UUID/PCI ID, driver and library versions), input shape/dtype and phase. The shared Triton operators are wired through it, and default and strict dispatch are unchanged.Because the kernels need to stay batch-invariant, candidates only vary launch parameters that keep the arithmetic order fixed. Every candidate is checked byte-for-byte against the fixed launch; all 101 kernel/config pairs match on B200. Tuning speeds up CUDA graph replay (up to about 30% on small elementwise ops) but not eager execution, so I've kept it off by default.
I haven't been able to validate on ROCm hardware, so that part of the issue is still open. If anyone with ROCm access would like to try it, I'd be happy to help.
Any feedback on the approach would be much appreciated. Thank you!
feat(platforms): persistent kernel auto-tuning in the HAL (RL-Align#50)