Skip to content

[Bugfix] Disable TMA on Blackwell GPUs to fix Triton autotuner OOM in fla/solve_trilfix: disable TMA on Blackwell (sm_12x) to prevent Triton autotuner OO… #36325

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
Rks2302 wants to merge 3 commits into
base: main
Choose a base branch
from
+5 −2
Open

[Bugfix] Disable TMA on Blackwell GPUs to fix Triton autotuner OOM in fla/solve_trilfix: disable TMA on Blackwell (sm_12x) to prevent Triton autotuner OO… #36325

Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
163f559
fix: disable TMA on Blackwell (sm_12x) to prevent Triton autotuner OO…
Rks2302 Mar 7, 2026
95d07ff
fix: use range check for TMA support to avoid redundant capability call
Rks2302 Mar 7, 2026
ae40230
fix: apply pre-commit formatting fixes
Rks2302 Mar 7, 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
7 changes: 5 additions & 2 deletions vllm/model_executor/layers/fla/ops/utils.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -152,10 +152,13 @@ def _check_platform() -> Literal["nvidia", "amd", "intel", "musa"]:
)
use_cuda_graph = is_nvidia and os.environ.get("FLA_USE_CUDA_GRAPH", "0") == "1"
is_gather_supported = hasattr(triton.language, "gather")
is_tma_supported = (is_nvidia and torch.cuda.get_device_capability(0)[0] >= 9) and (

is_tma_supported = (
is_nvidia and 9 <= torch.cuda.get_device_capability(0)[0] < 12
) and (
hasattr(triton.language, "_experimental_make_tensor_descriptor")
or hasattr(triton.language, "make_tensor_descriptor")
)
) # Upper bound < 12 disables TMA on Blackwell (sm_12x): Triton autotuner OOM


def get_all_max_shared_mem():
Expand Down