Skip to content

fix(gdn): use physical SM count for SM100 persistent prefill kernel #3155

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

Merged
kahyunnam merged 4 commits into from
Apr 25, 2026
+4 −4
Merged

fix(gdn): use physical SM count for SM100 persistent prefill kernel #3155

Changes from 1 commit
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
66e738c
fix(gdn): use physical SM count for SM100 persistent prefill kernel
arpera Apr 23, 2026
1f337a3
Update flashinfer/gdn_kernels/blackwell/gdn_prefill.py
arpera Apr 23, 2026
c7f1020
fix(gdn): drop get_max_active_clusters indirection in SM100 prefill
arpera Apr 23, 2026
92b4558
Merge branch 'main' into fix-backwell-gdn
arpera Apr 24, 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
8 changes: 4 additions & 4 deletions flashinfer/gdn_kernels/blackwell/gdn_prefill.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -33,9 +33,10 @@
import cuda.bindings.driver as cuda
import cutlass
import cutlass.cute as cute
import cutlass.utils as cutlass_utils
from cutlass.cute.runtime import from_dlpack

from flashinfer.cute_dsl.utils import get_max_active_clusters, get_num_sm
Comment thread
coderabbitai[bot] marked this conversation as resolved.
Outdated

from .gated_delta_net_chunked import GatedDeltaNetChunkedKernel


Expand Down Expand Up @@ -157,9 +158,8 @@ def chunk_gated_delta_rule_sm100(

if "compiled" not in cache:
# --- First call: compile the kernel ---
hardware_info = cutlass_utils.HardwareInfo()
num_sm = hardware_info.get_max_active_clusters(1)
max_active_clusters = hardware_info.get_max_active_clusters(1)
num_sm = get_num_sm(q.device)
max_active_clusters = min(get_max_active_clusters(1), num_sm)
Comment thread
arpera marked this conversation as resolved.
Outdated

gdn = GatedDeltaNetChunkedKernel(
io_dtype=io_dtype,
Expand Down
Loading