Skip to content

Commit 51ec5e4

Browse files
committed
[None][feat] Add SiTU support to Rubin FC12 MoE backend
Signed-off-by: Xin Guan <294044352+xguannv@users.noreply.github.com>
1 parent 0921d98 commit 51ec5e4

6 files changed

Lines changed: 466 additions & 130 deletions

File tree

‎tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py‎

Lines changed: 43 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -17773,7 +17773,7 @@ def _(
1777317773
# builds that do not ship cutlass.utils.rubin_helpers (SM107 only).
1777417774
if IS_CUTLASS_DSL_RUBIN_AVAILABLE:
1777517775
# ----------------------------------------------------------------
17776-
# Rubin NVFP4 Fused FC12 (FC1 gather+SwiGLU + FC2 finalize in ONE kernel)
17776+
# Rubin NVFP4 Fused FC12 (FC1 gather+gated-act + FC2 finalize in ONE kernel)
1777717777
# ----------------------------------------------------------------
1777817778
# Compat shim: the delivered fused kernel accesses ``cutlass.memory.*``
1777917779
# (SmemAllocator / TmemAllocator / get_smem_capacity_in_bytes) and
@@ -17866,16 +17866,20 @@ class Sm107BlockScaledContiguousGroupedGemmFusedFc12Runner(
1786617866
kernel_cache = dict()
1786717867
tuning_config_cache = dict()
1786817868

17869-
def __init__(self,
17870-
num_experts: int,
17871-
top_k: int,
17872-
num_local_experts: int,
17873-
local_expert_offset: int,
17874-
tile_size: int,
17875-
scaling_vector_size: int = 16,
17876-
swiglu_limit: float = float("inf"),
17877-
ep_size: int = 1,
17878-
enable_alltoall: bool = False):
17869+
def __init__(
17870+
self,
17871+
num_experts: int,
17872+
top_k: int,
17873+
num_local_experts: int,
17874+
local_expert_offset: int,
17875+
tile_size: int,
17876+
scaling_vector_size: int = 16,
17877+
swiglu_limit: float = float("inf"),
17878+
ep_size: int = 1,
17879+
enable_alltoall: bool = False,
17880+
activation_type: ActivationType = ActivationType.Swiglu,
17881+
situ_beta: Optional[float] = None,
17882+
situ_linear_beta: Optional[float] = None):
1787917883
super().__init__()
1788017884
self.num_experts = num_experts
1788117885
self.top_k = top_k
@@ -17884,6 +17888,9 @@ def __init__(self,
1788417888
self.tile_size = tile_size
1788517889
self.scaling_vector_size = scaling_vector_size
1788617890
self.swiglu_limit = swiglu_limit
17891+
self.activation_type = ActivationType(activation_type)
17892+
self.situ_beta = situ_beta
17893+
self.situ_linear_beta = situ_linear_beta
1788717894
# Used only by the in-op output memset (moved here so the memset
1788817895
# is the fused kernel's immediate stream predecessor).
1788917896
self.ep_size = ep_size
@@ -17908,6 +17915,9 @@ def unique_id(self):
1790817915
self.tile_size,
1790917916
self.scaling_vector_size,
1791017917
self.swiglu_limit,
17918+
int(self.activation_type),
17919+
self.situ_beta,
17920+
self.situ_linear_beta,
1791117921
)
1791217922

1791317923
def get_valid_tactics(
@@ -18205,7 +18215,8 @@ def forward(self, inputs: List[torch.Tensor],
1820518215
cache_key = (self.scaling_vector_size, self.tile_size,
1820618216
self.top_k, mma_tiler, mma_inst_shape,
1820718217
cluster_shape_mn, max_active_clusters,
18208-
self.swiglu_limit)
18218+
self.swiglu_limit, int(self.activation_type),
18219+
self.situ_beta, self.situ_linear_beta)
1820918220
if cache_key not in self.__class__.kernel_cache:
1821018221
gemm = self.__class__.kernel_class(
1821118222
self.scaling_vector_size,
@@ -18217,6 +18228,9 @@ def forward(self, inputs: List[torch.Tensor],
1821718228
use_pdl=True,
1821818229
swiglu_limit=self.swiglu_limit,
1821918230
scheduler="l2_atomic",
18231+
activation_type=self.activation_type,
18232+
situ_beta=self.situ_beta,
18233+
situ_linear_beta=self.situ_linear_beta,
1822018234
)
1822118235

1822218236
@cute.jit
@@ -18452,6 +18466,9 @@ def _run_nvfp4_fc12_fused_rubin(
1845218466
ep_size: int,
1845318467
enable_alltoall: bool,
1845418468
tuner_key: str,
18469+
activation_type: ActivationType = ActivationType.Swiglu,
18470+
situ_beta: Optional[float] = None,
18471+
situ_linear_beta: Optional[float] = None,
1845518472
) -> torch.Tensor:
1845618473
tuner = AutoTuner.get()
1845718474
runner = Sm107BlockScaledContiguousGroupedGemmFusedFc12Runner(
@@ -18464,6 +18481,9 @@ def _run_nvfp4_fc12_fused_rubin(
1846418481
swiglu_limit=swiglu_limit,
1846518482
ep_size=ep_size,
1846618483
enable_alltoall=enable_alltoall,
18484+
activation_type=activation_type,
18485+
situ_beta=situ_beta,
18486+
situ_linear_beta=situ_linear_beta,
1846718487
)
1846818488
# Input order matches Fc12FusedInputsHelper (FC1 prefix 0..9 mirrors
1846918489
# GatherGroupedGemmInputsHelper; FC2 tensors 10..14; expanded_idx 15
@@ -18500,7 +18520,9 @@ def _run_nvfp4_fc12_fused_rubin(
1850018520
"SymInt local_expert_offset, SymInt tile_size, float swiglu_limit, "
1850118521
"SymInt ep_size, bool enable_alltoall, "
1850218522
"SymInt scaling_vector_size=16, "
18503-
"str? precomputed_tactic=None) -> ()",
18523+
"str? precomputed_tactic=None, "
18524+
f"SymInt activation_type={int(ActivationType.Swiglu)}, "
18525+
"float? situ_beta=None, float? situ_linear_beta=None) -> ()",
1850418526
device_types="cuda")
1850518527
def cute_dsl_nvfp4_fc12_fused_rubin(
1850618528
input: torch.Tensor,
@@ -18529,6 +18551,9 @@ def cute_dsl_nvfp4_fc12_fused_rubin(
1852918551
enable_alltoall: bool,
1853018552
scaling_vector_size: int = 16,
1853118553
precomputed_tactic: Optional[str] = None,
18554+
activation_type: int = int(ActivationType.Swiglu),
18555+
situ_beta: Optional[float] = None,
18556+
situ_linear_beta: Optional[float] = None,
1853218557
) -> None:
1853318558
# In-place: finalize scatter-adds into ``output`` (mutates_args);
1853418559
# the op returns nothing so it does not alias its own input.
@@ -18541,7 +18566,8 @@ def cute_dsl_nvfp4_fc12_fused_rubin(
1854118566
local_expert_offset, tile_size, scaling_vector_size,
1854218567
swiglu_limit, precomputed_tactic, expanded_idx_to_permuted_idx,
1854318568
ep_size, enable_alltoall,
18544-
"trtllm::cute_dsl_nvfp4_fc12_fused_rubin")
18569+
"trtllm::cute_dsl_nvfp4_fc12_fused_rubin",
18570+
ActivationType(activation_type), situ_beta, situ_linear_beta)
1854518571

1854618572
@torch.library.register_fake("trtllm::cute_dsl_nvfp4_fc12_fused_rubin")
1854718573
def _(
@@ -18571,5 +18597,8 @@ def _(
1857118597
enable_alltoall: bool,
1857218598
scaling_vector_size: int = 16,
1857318599
precomputed_tactic: Optional[str] = None,
18600+
activation_type: int = int(ActivationType.Swiglu),
18601+
situ_beta: Optional[float] = None,
18602+
situ_linear_beta: Optional[float] = None,
1857418603
) -> None:
1857518604
return None

0 commit comments

Comments
 (0)