[CUDA] Enable tanh-GEGLU for MoE and integer QMoE

docs/ContribOperators.md

@@ -4597,6 +4597,10 @@ This version of the operator has been available since version 1 of the 'com.micr
         When swiglu_fusion=0, two GEMMs are not fused, and they are FC1 and FC3 in the inputs.
         When swiglu_fusion=1, two GEMMs are fused so that g and l are computed in a single GEMM (FC1), and g and l are interleaved on each row of size 2 * inter_size.
         When swiglu_fusion=2, two GEMMs are fused, and g and l are concatenated on each row.
+        CUDA also supports GEGLU: geglu = GELU(g, approximate='tanh') * l.
+        GEGLU requires fused FC1 projections: swiglu_fusion=1 stores [g0,l0,g1,l1,...],
+        and swiglu_fusion=2 stores [g | l]. Separate FC3 inputs are not supported.
+        activation_alpha, activation_beta and swiglu_limit do not affect GEGLU.
 
 
 #### Version
@@ -4611,15 +4615,15 @@ This version of the operator has been available since version 1 of the 'com.micr
 <dt><tt>activation_beta</tt> : float</dt>
 <dd>Beta parameter used in activation function.</dd>
 <dt><tt>activation_type</tt> : string</dt>
-<dd>Activation function to use. Choose from relu, gelu, silu, swiglu and identity. CUDA also supports relu2 (squared ReLU). Default is relu</dd>
+<dd>Activation function to use. Choose from relu, gelu, silu, swiglu and identity. CUDA also supports relu2 (squared ReLU) and geglu (GELU-tanh(gate) * up). GEGLU requires fused FC1 gate/up weights, swiglu_fusion=1 or 2, and no FC3 inputs. Default is relu</dd>
 <dt><tt>k</tt> : int</dt>
 <dd>Number of top experts to select from expert pool</dd>
 <dt><tt>normalize_routing_weights</tt> : int</dt>
 <dd>Whether to normalize routing weights</dd>
 <dt><tt>swiglu_fusion</tt> : int</dt>
-<dd>0: not fused, 1: fused and interleaved. 2: fused and not interleaved.</dd>
+<dd>Gate/up layout for swiglu or geglu. 0: not fused (swiglu only), 1: fused and interleaved, 2: fused and not interleaved ([gate | up] for geglu).</dd>
 <dt><tt>swiglu_limit</tt> : float</dt>
-<dd>The limit used to clamp in SwiGLU. No clamp when limit is not provided.</dd>
+<dd>The limit used to clamp in SwiGLU. No clamp when limit is not provided. Does not affect GEGLU.</dd>
 <dt><tt>use_sparse_mixer</tt> : int</dt>
 <dd>Whether to use sparse mixer</dd>
 </dl>
@@ -4632,7 +4636,7 @@ This version of the operator has been available since version 1 of the 'com.micr
 <dt><tt>router_probs</tt> : T</dt>
 <dd>2D input tensor with shape (total_tokens, num_experts), where total_tokens must match the flattened token count of input</dd>
 <dt><tt>fc1_experts_weights</tt> : T</dt>
-<dd>3D input tensor with shape (num_experts, fusion_size * inter_size, hidden_size), where fusion_size is 2 for fused swiglu, and 1 otherwise</dd>
+<dd>3D input tensor with shape (num_experts, fusion_size * inter_size, hidden_size), where fusion_size is 2 for fused swiglu or geglu, and 1 otherwise</dd>
 <dt><tt>fc1_experts_bias</tt> (optional) : T</dt>
 <dd>2D optional input tensor with shape (num_experts, fusion_size * inter_size)</dd>
 <dt><tt>fc2_experts_weights</tt> : T</dt>
@@ -6452,6 +6456,10 @@ This version of the operator has been available since version 1 of the 'com.micr
         When swiglu_fusion=0, two GEMMs are not fused, and they are FC1 and FC3 in the inputs.
         When swiglu_fusion=1, two GEMMs are fused so that g and l are computed in a single GEMM (FC1), and g and l are interleaved on each row of size 2 * inter_size.
         When swiglu_fusion=2, two GEMMs are fused, and g and l are concatenated on each row.
+        CUDA also supports GEGLU with quant_type='int': geglu = GELU(g, approximate='tanh') * l.
+        GEGLU requires fused FC1 projections: swiglu_fusion=1 stores [g0,l0,g1,l1,...],
+        and swiglu_fusion=2 stores [g | l]. Separate FC3 weights are not supported.
+        activation_alpha, activation_beta and swiglu_limit do not affect GEGLU.
 
     The raw INT3 profile extends version 1; execution support is provider-specific.
     For INT3, the requirements below supersede the legacy format descriptions above.
@@ -6483,7 +6491,7 @@ This version of the operator has been available since version 1 of the 'com.micr
 <dt><tt>activation_beta</tt> : float</dt>
 <dd>Beta parameter used in activation function.</dd>
 <dt><tt>activation_type</tt> : string</dt>
-<dd>Activation function to use. Choose from relu, gelu, silu, swiglu and identity. CUDA also supports relu2 (squared ReLU). Default is relu</dd>
+<dd>Activation function to use. Choose from relu, gelu, silu, swiglu and identity. CUDA also supports relu2 (squared ReLU) and geglu (GELU-tanh(gate) * up). GEGLU requires quant_type='int', fused FC1 gate/up weights, swiglu_fusion=1 or 2, and no separate FC3 weights. Default is relu</dd>
 <dt><tt>block_size</tt> : int</dt>
 <dd>For integer quantization, size of each quantization block along the K (input feature) dimension. Must be power of two and ≥ 16 (e.g., 16, 32, 64, 128). For integer and FP4 quantization, both hidden_size and inter_size must be divisible by the block size. FP8 with block_size=128 supports partial 128x128 tiles. The FP4 modes always use blocking: MXFP4 ('fp4'/'wfp4afp8') is normalized to block_size 32 and NVFP4 ('nvfp4') to block_size 16, even when block_size is omitted. For FP8 ('fp8'), a positive value instead specifies square blocks along both N and K, with floating-point scales shaped [E, ceil(N/block_size), ceil(K/block_size)]; partial blocks are allowed. Without a positive value, FP8 uses per-expert global scales. For integer quantization ('int'), omitting block_size means there is no blocking and a whole column shares one scaling factor. </dd>
 <dt><tt>expert_weight_bits</tt> : int</dt>
@@ -6501,9 +6509,9 @@ This version of the operator has been available since version 1 of the 'com.micr
 <dt><tt>quant_type</tt> : string</dt>
 <dd>Quantization type: 'int' for integer quantization (default), 'fp4' for MXFP4 quantization, 'nvfp4' for NVFP4 quantization, 'fp8' for FP8 e4m3 weight-only quantization, or 'wfp4afp8' for MXFP4 weight with FP8 activation. When quant_type is 'fp4' or 'nvfp4', weights are stored in E2M1 FP4 format (2 values per byte), fc*_scales inputs contain the FP4 block scales, and fc*_global_scale inputs must be provided. 'fp4' uses Float8E8M0 block scales with block_size 32; 'nvfp4' uses Float8E4M3FN block scales with block_size 16. 'fp8' uses float8e4m3fn weights and, when block_size > 0, float32, float16, or bfloat16 square-block scales in fc*_scales instead of fc*_global_scale. The INT3 profile requires integer quantization.</dd>
 <dt><tt>swiglu_fusion</tt> : int</dt>
-<dd>0: not fused, 1: fused and interleaved. 2: fused and not interleaved.</dd>
+<dd>Gate/up layout for swiglu or geglu. 0: not fused (swiglu only), 1: fused and interleaved, 2: fused and not interleaved ([gate | up] for geglu).</dd>
 <dt><tt>swiglu_limit</tt> : float</dt>
-<dd>The limit used to clamp inputs in SwiGLU. It is infinite when limit is not provided.</dd>
+<dd>The limit used to clamp inputs in SwiGLU. It is infinite when limit is not provided. Does not affect GEGLU.</dd>
 <dt><tt>use_sparse_mixer</tt> : int</dt>
 <dd>Whether to use sparse mixer</dd>
 <dt><tt>weights_prepacked</tt> : int</dt>
@@ -6518,7 +6526,7 @@ This version of the operator has been available since version 1 of the 'com.micr
 <dt><tt>router_probs</tt> : T</dt>
 <dd>2D tensor with shape (total_tokens, num_experts), where total_tokens must match the flattened token count of input</dd>
 <dt><tt>fc1_experts_weights</tt> : T1</dt>
-<dd>3D tensor with shape (num_experts, fusion_size * inter_size, hidden_size * effective_fc1_bits / 8). The last dimension must be byte-aligned. The fusion_size is 2 for fused swiglu, or 1 otherwise. effective_fc1_bits is fc1_expert_weight_bits when provided, otherwise expert_weight_bits.</dd>
+<dd>3D tensor with shape (num_experts, fusion_size * inter_size, hidden_size * effective_fc1_bits / 8). The last dimension must be byte-aligned. The fusion_size is 2 for fused swiglu or geglu, or 1 otherwise. effective_fc1_bits is fc1_expert_weight_bits when provided, otherwise expert_weight_bits.</dd>
 <dt><tt>fc1_scales</tt> (optional) : T2</dt>
 <dd>Optional weight scales. For quant_type='int', this is a 2D tensor with shape (num_experts, fusion_size * inter_size), or a 3D tensor with shape (num_experts, fusion_size * inter_size, hidden_size / block_size) when block_size is provided. For quant_type='fp4' or 'wfp4afp8', this is a float8e8m0 MXFP block-scale tensor with shape (num_experts, fusion_size * inter_size, hidden_size / 32). For quant_type='nvfp4', this is a float8e4m3fn NVFP4 block-scale tensor with shape (num_experts, fusion_size * inter_size, hidden_size / 16). For quant_type='fp8' and block_size > 0, required float32, float16, or bfloat16 scales with shape (num_experts, ceil(fusion_size * inter_size / block_size), ceil(hidden_size / block_size)).</dd>
 <dt><tt>fc1_experts_bias</tt> (optional) : T</dt>

docs/contrib_ops/cuda/moe_qmoe.md

@@ -59,7 +59,7 @@ The execution pipeline is:
 
 ```
 input tokens → router (top-k softmax) → permute by expert
-  → GEMM1 (per-expert) → activation (SiLU/GeLU/ReLU/SwiGLU)
+  → GEMM1 (per-expert) → activation (SiLU/GeLU/ReLU/SwiGLU/GEGLU)
   → GEMM2 (per-expert) → un-permute → weighted sum
 ```
 
@@ -72,7 +72,7 @@ input tokens → router (top-k softmax) → permute by expert
 | Attribute | Type | Default | Description |
 |-----------|------|---------|-------------|
 | `k` | int | 1 | Top-K experts selected per token. |
-| `activation_type` | string | `"relu"` | One of `"relu"`, `"gelu"`, `"silu"`, `"swiglu"`, `"identity"`. These are the only values accepted end-to-end (attribute parsing); other kernel-internal types are not reachable from ONNX. |
+| `activation_type` | string | `"relu"` | One of `"relu"`, `"relu2"`, `"gelu"`, `"silu"`, `"swiglu"`, `"geglu"`, `"identity"`. CUDA GEGLU computes `GELU(gate, approximate="tanh") * up`, requires fused FC1 and `swiglu_fusion=1` or `2`; QMoE additionally requires `quant_type="int"`. |
 | `normalize_routing_weights` | int | 0 | Re-normalize the top-k weights to sum to 1. |
 | `use_sparse_mixer` | int | 0 | Enable sparse-mixer routing variant. |
 | `swiglu_fusion` | int | 0 | 0=no fusion, 1=interleaved (Gate/Value), 2=block (Gate;Value). See [§8](#8-swiglu-fusion). |
@@ -123,12 +123,13 @@ to the selected `quant_type` are simply omitted (most are `Optional`).
 | 20 | `fc2_act_block_scale` | T2 (Opt, float8e8m0) | `(E, M_pad, inter/32)` | wfp4afp8 (Variant B) |
 
 `E = num_experts`. `pack = 8 / expert_weight_bits` for INT/MXFP4 weights; `pack = 1`
-for FP8 weights. `fusion = 2` for `swiglu_fusion=1`, otherwise `1`.
+for FP8 weights. `fusion = 2` for fused SwiGLU or GEGLU with
+`swiglu_fusion=1` or `2`, otherwise `1`.
 
 > **CUDA EP restriction for `quant_type="int"`:** `fc1_scales` and `fc2_scales`
 > must use the same FP16 or BF16 type as the activation input. Although the
 > operator schema also permits float scales, the CUDA INT kernel does not
-> currently cast float scales internally. `fusion = 2` for `swiglu_fusion=1`, otherwise `1`.
+> currently cast float scales internally.
 
 For packed input, tokens from variable-length sequences may be concatenated
 without padding. MoE routing is token-local, so no sequence-offset input is
@@ -636,6 +637,26 @@ easy to feed Mixtral-style models without offline fusion.
 
 ---
 
+### GEGLU
+
+CUDA `MoE` and integer `QMoE` also accept `activation_type="geglu"`.
+They evaluate `GELU(gate, approximate="tanh") * up`, not a scaled-SwiGLU
+approximation or erf-based GELU. The fused FC1 output has `2 * inter_size`
+columns: fusion mode 1 stores `[gate0, up0, gate1, up1, ...]`, and mode 2
+stores `[gate | up]`. GEGLU requires one of these explicit fusion modes and
+does not support separate FC3 weights. `activation_alpha`, `activation_beta`,
+and `swiglu_limit` apply to SwiGLU only.
+
+GEGLU uses the generic gated activation path; specialized kernels that
+hardcode SwiGLU remain excluded. Integer INT4/INT8 still use weight-only
+CUTLASS runners. GEGLU is rejected for FP4, NVFP4, FP8, and WFP4AFP8 rather
+than silently selecting a different activation. CPU and other execution
+providers do not gain GEGLU support from this CUDA implementation.
+
+The activation and product are computed in float before conversion to the
+activation output type. This is not a promise of bitwise parity with eager
+PyTorch code that materializes GELU in FP16/BF16 before multiplication.
+
 ## 9. FP4 (MXFP4) Details
 
 The QMoE operator supports **MXFP4** quantized weights with FP16/BF16 activations

onnxruntime/contrib_ops/cuda/llm/moe_gemm/moe_gemm_activation_kernels.cuh

@@ -131,7 +131,7 @@ void doGatedActivation(ActivationOutputType* output, const GemmOutputType* gemm_
       case ActivationType::SwigluBias:
         return &doGatedActivationKernel<ActivationOutputType, GemmOutputType, SiLu>;
       case ActivationType::Geglu:
-        return &doGatedActivationKernel<ActivationOutputType, GemmOutputType, GELU>;
+        return &doGatedActivationKernel<ActivationOutputType, GemmOutputType, GELU_taylor>;
       case ActivationType::Silu:
         return &doGatedActivationKernel<ActivationOutputType, GemmOutputType, SiLu>;
       case ActivationType::Gelu:
@@ -226,6 +226,10 @@ __global__ void doActivationKernel(T* output, const GemmOutputType* gemm_result,
     const int64_t num_elems_in_col = inter_size / ACTIVATION_ELEM_PER_THREAD;
     assert(gated_off % ACTIVATION_ELEM_PER_THREAD == 0);
     const int64_t gated_off_vec = gated_off / ACTIVATION_ELEM_PER_THREAD;
+    // Chunked GEGLU is [gate | up]; preserve the legacy SwiGLU chunk order.
+    const int64_t gate_offset =
+        gated && activation_params.activation_type == ActivationType::Geglu ? 0 : gated_off_vec;
+    const int64_t linear_offset = gated_off_vec - gate_offset;
 
     ActFn<ComputeElem> fn{};
     constexpr bool IsSiLu = std::is_same_v<ActFn<float>, cutlass::epilogue::thread::SiLu<float>>;
@@ -244,11 +248,9 @@ __global__ void doActivationKernel(T* output, const GemmOutputType* gemm_result,
           linear_part[i] = static_cast<float>(scalar_gemm[2 * global_elem + 1]);
         }
       } else {
-        // If not gated, gate_part reads from elem_index (gated_off_vec is 0).
-        // If gated, gate_part reads from elem_index + inter_size_vec (chunk 1).
-        gate_part = arrayConvert<GemmResultElem, ComputeElem>(gemm_result_vec[elem_index + gated_off_vec]);
+        gate_part = arrayConvert<GemmResultElem, ComputeElem>(gemm_result_vec[elem_index + gate_offset]);
         if (gated) {
-          linear_part = arrayConvert<GemmResultElem, ComputeElem>(gemm_result_vec[elem_index]);
+          linear_part = arrayConvert<GemmResultElem, ComputeElem>(gemm_result_vec[elem_index + linear_offset]);
         }
       }
 
@@ -261,9 +263,9 @@ __global__ void doActivationKernel(T* output, const GemmOutputType* gemm_result,
             linear_part[i] += static_cast<float>(scalar_bias[2 * global_elem + 1]);
           }
         } else {
-          gate_part = gate_part + arrayConvert<BiasElem, ComputeElem>(bias_ptr_vec[elem_index + gated_off_vec]);
+          gate_part = gate_part + arrayConvert<BiasElem, ComputeElem>(bias_ptr_vec[elem_index + gate_offset]);
           if (gated) {
-            linear_part = linear_part + arrayConvert<BiasElem, ComputeElem>(bias_ptr_vec[elem_index]);
+            linear_part = linear_part + arrayConvert<BiasElem, ComputeElem>(bias_ptr_vec[elem_index + linear_offset]);
           }
         }
       }
@@ -390,8 +392,9 @@ void doActivation(T* output, const GemmOutputType* gemm_result, const float* fp8
         case ActivationType::Identity:
           return &doActivationKernel<T, GemmOutputType, ScaleBiasType, Identity, decltype(block_scaling_type)::value>;
         case ActivationType::Gelu:
-        case ActivationType::Geglu:
           return &doActivationKernel<T, GemmOutputType, ScaleBiasType, GELU, decltype(block_scaling_type)::value>;
+        case ActivationType::Geglu:
+          return &doActivationKernel<T, GemmOutputType, ScaleBiasType, GELU_taylor, decltype(block_scaling_type)::value>;
         case ActivationType::Relu:
           return &doActivationKernel<T, GemmOutputType, ScaleBiasType, ReLu, decltype(block_scaling_type)::value>;
         case ActivationType::Relu2:

onnxruntime/contrib_ops/cuda/moe/moe.cc

@@ -72,9 +72,13 @@ Status MoE<T>::ComputeInternal(OpKernelContext* context) const {
     LogSwigluFusionRemapOnce();
   }
 
-  bool is_fused_swiglu = (activation_type_ == ActivationType::Swiglu) &&
-                         (swiglu_fusion != 0) &&
-                         (fc3_experts_weights_optional == nullptr);
+  ORT_RETURN_IF(activation_type_ == ActivationType::Geglu &&
+                    (fc3_experts_weights_optional != nullptr || fc3_experts_bias_optional != nullptr),
+                "GEGLU requires gate/up projections fused into FC1, without FC3 inputs.");
+  bool is_fused_gated = (activation_type_ == ActivationType::Swiglu ||
+                         activation_type_ == ActivationType::Geglu) &&
+                        (swiglu_fusion != 0) &&
+                        (fc3_experts_weights_optional == nullptr);
 
   MoEParameters moe_params;
   ORT_RETURN_IF_ERROR(::onnxruntime::contrib::moe_helper::CheckInputs<Tensor>(
@@ -83,7 +87,7 @@ Status MoE<T>::ComputeInternal(OpKernelContext* context) const {
       fc2_experts_weights, fc2_experts_bias_optional, nullptr, nullptr,
       fc3_experts_weights_optional, fc3_experts_bias_optional, nullptr, nullptr,
       1,  //  no quantization so pack size is 1
-      is_fused_swiglu,
+      is_fused_gated,
       0));  // no block-wise quantization for regular MoE
   ORT_RETURN_IF_NOT(k_ > 0 && k_ <= moe_params.num_experts,
                     "MoE requires 0 < k <= num_experts, got k=", k_,

onnxruntime/contrib_ops/cuda/moe/moe_base.h

@@ -57,6 +57,8 @@ class MoEBase {
       activation_type_ = ActivationType::Silu;
     } else if (activation_type_str == "swiglu") {
       activation_type_ = ActivationType::Swiglu;
+    } else if (activation_type_str == "geglu") {
+      activation_type_ = ActivationType::Geglu;
     } else if (activation_type_str == "identity") {
       activation_type_ = ActivationType::Identity;
     } else {
@@ -81,8 +83,11 @@ class MoEBase {
     swiglu_fusion_ = static_cast<int>(op_kernel_info.GetAttrOrDefault<int64_t>("swiglu_fusion", 0));
     ORT_ENFORCE(swiglu_fusion_ >= 0 && swiglu_fusion_ <= 2,
                 "swiglu_fusion must be 0, 1, or 2, but got ", swiglu_fusion_);
-    ORT_ENFORCE(activation_type_ == ActivationType::Swiglu || swiglu_fusion_ == 0,
-                "swiglu_fusion is only valid when activation_type is 'swiglu'.");
+    ORT_ENFORCE(activation_type_ == ActivationType::Swiglu ||
+                    activation_type_ == ActivationType::Geglu || swiglu_fusion_ == 0,
+                "swiglu_fusion is only valid when activation_type is 'swiglu' or 'geglu'.");
+    ORT_ENFORCE(activation_type_ != ActivationType::Geglu || swiglu_fusion_ != 0,
+                "GEGLU requires fused FC1 gate/up weights and swiglu_fusion=1 or 2.");
 
     // SwiGLU limit for clamping (optional, use infinity if not provided)
     swiglu_limit_ = op_kernel_info.GetAttrOrDefault<float>("swiglu_limit", std::numeric_limits<float>::infinity());