From fca72427ea8c2bbb9e52fe5d9f763f7288efc431 Mon Sep 17 00:00:00 2001 From: biondizzle Date: Tue, 2 Jun 2026 08:11:18 +0000 Subject: [PATCH] fix: add fp4_out/sf_out/l2_global_scale params to fused_swiglu kernel() signature The __call__ method passes these 3 Optional params to self.kernel(), but kernel() didn't accept them, causing TypeError: too many positional arguments during cute.compile(). This was the CuTeDSL 'arg-binding bug' blocking P0/P1. --- dsv4/kernels/gemm/fused_swiglu.py | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/dsv4/kernels/gemm/fused_swiglu.py b/dsv4/kernels/gemm/fused_swiglu.py index 4040666c..61fcffeb 100644 --- a/dsv4/kernels/gemm/fused_swiglu.py +++ b/dsv4/kernels/gemm/fused_swiglu.py @@ -1285,6 +1285,10 @@ class FusedSwiGLUScaledGroupedGemmKernel: # ── Optional: NVFP4 per-expert global scales ── global_scale_a: Optional[cute.Tensor], global_scale_b: Optional[cute.Tensor], + # ── Fused SwiGLU epilogue outputs (replaces out when fused_swiglu=True) ── + fp4_out: Optional[cute.Tensor] = None, + sf_out: Optional[cute.Tensor] = None, + l2_global_scale: Optional[cute.Tensor] = None, ): """ GPU device kernel for MoE Scaled Grouped GEMM with block scaling.