Autotuner NoConfigFound on RTX 5090 (Blackwell sm_120) — manual configs work fine
## Description
Helion's autotuner fails with `NoConfigFound` on NVIDIA RTX 5090 (Blackwell, compute capability 12.0 / sm_120). However, manually specifying a config works perfectly — the kernel compiles, runs, and produces correct results.
This suggests the autotuner is generating configs with parameters that are invalid on Blackwell, while the underlying Triton compilation and execution work fine on this architecture.
## Reproduction
**Fails (autotuner):**
```python
import torch
import helion
import helion.language as hl
@helion.kernel()
def vector_add(a: torch.Tensor, b: torch.Tensor) -> torch.Tensor:
n, = a.size()
out = torch.empty_like(a)
for tile_n in hl.tile(n):
out[tile_n] = a[tile_n] + b[tile_n]
return out
a = torch.randn(1024, device="cuda")
b = torch.randn(1024, device="cuda")
result = vector_add(a, b) # NoConfigFound after ~200 configs tried
```
**Works (manual config):**
```python
from helion.runtime.config import Config
@helion.kernel(configs=[Config(block_sizes=[128])])
def vector_add(a: torch.Tensor, b: torch.Tensor) -> torch.Tensor:
n, = a.size()
out = torch.empty_like(a)
for tile_n in hl.tile(n):
out[tile_n] = a[tile_n] + b[tile_n]
return out
a = torch.randn(1024, device="cuda")
b = torch.randn(1024, device="cuda")
result = vector_add(a, b)
print(torch.allclose(result, a + b)) # True, zero error
```
## Autotuner behavior
With debug logging enabled:
- Initial population precompiles 100 configs at ~65 configs/s (compilation succeeds)
- All 200 configs fail at runtime
- `NoConfigFound` is raised
This suggests the generated Triton kernels compile to valid PTX but crash or produce incorrect results when executed on sm_120. A bare Triton kernel with the same logic (`@triton.jit` with manual block size) runs correctly on the same GPU.
## Environment
- **GPU:** NVIDIA GeForce RTX 5090 (Compute capability 12.0 / sm_120)
- **Helion:** 1.2.0
- **Triton:** 3.6.0
- **PyTorch:** 2.10.0+cu128
- **CUDA:** 12.8
- **OS:** Fedora 44 (kernel 6.19.10)
- **PyTorch arch list includes sm_120:** confirmed via `torch.cuda.get_arch_list()`
## Notes
- Bare Triton kernels (`@triton.jit`) work correctly on this GPU with manually specified block sizes
- The issue is isolated to the autotuner's config search, not the compilation or runtime pipeline
- `Config(block_sizes=[128])` works, so the autotuner may be exploring configs with warp counts, stage counts, or other parameters that are invalid on Blackwell
0 条评论