Skip to content

[FIX] Run no blocks for an empty grid instead of failing - #523

Open
mark14wu wants to merge 1 commit into
mainfrom
fix/empty-grid-launch
Open

mark14wu wants to merge 1 commit into
mainfrom
fix/empty-grid-launch

Conversation

@mark14wu

@mark14wu mark14wu commented Nov 1, 2026

Copy link
Copy Markdown
Collaborator

Summary

A launch whose grid has a zero dimension fails under every client. For example, a test that runs a kernel on an empty tensor with grid = (triton.cdiv(n, BLOCK),) and n == 0:

ValueError: max_workers must be greater than 0

_prepare_grid_launch sets max_workers = min(cfg.num_sms, total_blocks), which is 0 for an empty grid. _run_grid runs blocks sequentially only when max_workers == 1, so an empty grid took the thread-pool path, and ThreadPoolExecutor(max_workers=0) raised. Triton runs no blocks for an empty grid, natively and under TRITON_INTERPRET=1. _run_grid now returns before running any block when max_workers == 0.

Tests

  • New in tests/unit/test_multithreading.py: test_empty_grid_runs_no_blocks, a launch with grid (0,) under the tracer, the sanitizer, and the race detector, each with 1 and 2 emulated SMs. It checks that the launch completes and stores nothing. All 6 cases fail on main with the error above.
  • pytest tests/: 366 passed, 6 skipped on main; 372 passed, 6 skipped with this PR (Python 3.12, Triton 3.8.0).
  • ruff check and ruff format --check (v0.1.9) pass on the changed files.

🤖 Generated with Claude Code

https://claude.ai/code/session_016HMpCLXdUVNfMbJTj8aHE2

A launch whose grid has a zero dimension, such as
`grid = (triton.cdiv(n, BLOCK),)` with `n == 0`, set `max_workers` to
`min(num_sms, 0) == 0`. `_run_grid` then took the thread-pool path, and
`ThreadPoolExecutor(max_workers=0)` raised "max_workers must be greater
than 0", so every client failed on such a launch. Triton runs no blocks
for an empty grid, natively and under `TRITON_INTERPRET=1`; `_run_grid`
now returns without running any.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_016HMpCLXdUVNfMbJTj8aHE2
@github-actions

github-actions Bot commented Nov 1, 2026

Copy link
Copy Markdown

Performance Benchmark

Benchmark main (min) PR (min) Change Samples
gemm N/A N/A N/A 0 / 0
gemm_oob N/A N/A N/A 0 / 0
indirect_load N/A N/A N/A 0 / 0
nested_loop N/A N/A N/A 0 / 0
block_pointer_loop_advance N/A N/A N/A 0 / 0
liger_jsd N/A N/A N/A 0 / 0
flaggems_layernorm N/A N/A N/A 0 / 0
swiglu 0.119s 0.120s +0.8% 20 / 20
cross_entropy N/A N/A N/A 0 / 0
fused_linear_jsd N/A N/A N/A 0 / 0
Total N/A N/A N/A N/A

Iterations: 1 warmup + 20 measured
Samples are shown as main / PR; long pytest benchmarks may use fewer samples.

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant