[Common][PyTorch] Add SiTU-GLU activation - #3402
Conversation
Signed-off-by: Harry Zhou <hhanyu@nvidia.com>
Greptile SummaryThe PR adds native and row-scaled SiTU-GLU operations and conditionally integrates them with the fused grouped-MLP path.
Confidence Score: 5/5The PR appears safe to merge. No blocking failure remains; the previously reported recomputation no-op is resolved by rejecting unsupported scaled-GLU requests while leaving supported ScaledSReLU recomputation intact. Important Files Changed
Flowchart%%{init: {'theme': 'neutral'}}%%
flowchart LR
A[GroupedLinear FC1] --> B{cuDNN frontend supports SiTU-GLU parameters?}
B -->|Yes| C[Fused grouped SiTU-GLU path]
B -->|No| D[Native ScaledSiTUGLU fallback]
C --> E[GroupedLinear FC2]
D --> E
Reviews (8): Last reviewed commit: "[Common] Fix scaled activation scale-gra..." | Re-trigger Greptile |
Signed-off-by: Harry Zhou <hhanyu@nvidia.com>
Signed-off-by: Harry Zhou <hhanyu@nvidia.com>
Signed-off-by: Harry Zhou <hhanyu@nvidia.com>
|
Update in TE now feature-detects The test update also makes fused-path validation strict: it traces the real cuDNN wrappers, verifies forward/backward activation names and generated tensor dtypes, and does not accept silent native fallback when fusion is required. Validation on NVIDIA B200:
|
|
/te-ci |
Signed-off-by: Harry Zhou <hhanyu@nvidia.com>
Signed-off-by: Harry Zhou <hhanyu@nvidia.com>
|
Addressed the gated-derivative review feedback in 15dd10f. The native FP8, MXFP8 colwise/rowwise, vectorized, and row-scaled implementations now use one contract:
For ordinary GLUs, dgate_elt is 1. For clamped GLUs it is the 0/1 clamp derivative, and for SiTU-GLU it is the smooth SiTU up-branch derivative. This removes the SiTU-only conditional composition and makes the row-scaled helper follow the same semantics as the FP8/MXFP8 kernels. Focused validation on B300 with PyTorch 26.06, cuDNN frontend 1.26.0 source, and CuTe DSL 4.6.2:
Related cuDNN-frontend patch for the SiTU-GLU Hadamard API contracts: NVIDIA/cudnn-frontend#670 |
Co-authored-by: Tim Moon <4406448+timmoon10@users.noreply.github.com> Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com>
Signed-off-by: Harry Zhou <hhanyu@nvidia.com>
|
/te-ci |
| } | ||
| gate_elt = min(max(-p.limit, gate_elt), p.limit) + p.glu_linear_offset; | ||
| } else if constexpr (std::is_same<ParamOP, SiTUGLUParam>::value) { | ||
| dgate_elt = dsitu_up<float, float>(gate_elt, p); |
There was a problem hiding this comment.
One of the optimizations that we did for swiglu and clamped_swiglu was to do the exponential computations shared by both the activation and dactivation functions once, as you can see in this gated kernel.
For example in this case, situ and dsitu can share the tanh computation being done once and use it. Ideally compiler should be smart enough to do those shared computations once. But we did find it wasnt the case for swiglu and clamped_swiglu in terms of reusing exp computations.
Not a blocker for this PR, but we should definitely see if the perf of this kernel becomes better if do the same optimization
cc: @ptrendx @Oleg-Goncharov
The SiTU-GLU merge (NVIDIA#3402) brought _cudnn_frontend_supports_grouped_gemm_situglu() into this file, which asks inspect.signature(wrapper).parameters for the arguments it needs rather than comparing frontend versions -- the same conclusion this branch reached independently, now the house style. Two things to match. Guard the signature call with `except (TypeError, ValueError)`: a callable that is not introspectable answers "no" instead of raising out of a backward pass. I had left this out on the grounds that the wrapper is a plain undecorated function, which is true today but is not a property this code controls. And say "feature-detect" in the docstring summary, as the neighbor does. Also dropped the sentence about use_single_group_runtime_offsets from the docstring. The neighbor now demonstrates the pattern in the same file, so the cautionary tale is no longer what makes the choice legible. `import inspect` came in with the merge, so this branch no longer adds it. Signed-off-by: Zhiyu Li <zhiyul@nvidia.com>
Description
Add native and row-scaled SiTU-GLU operations and connect them to the fused grouped-MLP path when the installed cuDNN frontend exposes SiTU-GLU support.
SiTU-GLU is the soft-capped gated activation used by Kimi K3. For gate input
a, up inputb, and positive cap parametersbeta1,beta2, this PR computesThe backward kernels use
The default values
beta1=4andbeta2=25follow the Kimi K3 technical report.The grouped-GEMM integration depends on the cuDNN frontend support proposed at NVIDIA/cudnn-frontend#645. Compatibility is detected from the forward and backward wrapper signatures instead of from a package-version threshold. An older or unpatched cuDNN frontend therefore keeps the correct native
GroupedLinear -> ScaledSiTUGLU -> GroupedLinearfallback, while a frontend that exposes bothsitu_beta1andsitu_beta2parameters enables the fused path.Type of change
Changes
te.ops.SiTUGLUwith contiguous and block-interleaved layouts and existing FP8/MXFP8 quantizer integration.te.ops.ScaledSiTUGLU, including optional scale gradients and a constructor-compatible recomputation flag that explicitly rejects unsupportedTruerequests.activation_recompute_in_mlp=Truefor scaled SwiGLU, SiTU-GLU, and clamped QGeGLU instead of silently ignoring it in the fused grouped MLP. Scaled SReLU remains supported.ScaledSiTUGLUto grouped-MLP dimension validation and passsituglu/dsitugluplus both cap parameters to cuDNN frontend forward/backward wrappers.Performance and compatibility
nvidia-cudnn-frontend==1.26.0,nvidia-cutlass-dsl==4.5.0) correctly reports the feature as unavailable and uses the native fallback.Validation
Run on one NVIDIA B300 GPU in the 26.06 container using an isolated private venv:
NVTE_CUDA_ARCHS="100;103a;"passed. Optional NCCL-EP was disabled because the separately reused NCCL-EP headers predate upstream TE's current API; no SiTU code depends on NCCL-EP.24 passed: focused native/scaled SiTU-GLU tests with stock cuDNN frontend, including FP32/FP16/BF16, interleaving, invalid parameters, MXFP8 output/input-gradient quantization, scale gradients, unsupported recomputation rejection, feature detection, and grouped fallback.2 passed: grouped MXFP8 fallback with(beta1, beta2)=(4,25)and(2,8)on stock cuDNN frontend 1.26.0.2 passed: fused grouped MXFP8 forward/backward with the same two parameter pairs using [CuTeDSL] Add grouped SiTU-GLU activation cudnn-frontend#645.244 passed, 206 skipped: existing SwiGLU-focused native, quantized, grouped-MLP, and MCore-integration regression selection.Checklist