From 8213668c47b6e7dcd3a7a7b34b1e315614a21acf Mon Sep 17 00:00:00 2001
From: shihaobai <1798930569@qq.com>
Date: Thu, 30 Jul 2026 05:06:29 +0000
Subject: [PATCH 1/9] fix qwen3.5 gate rounding for RL consistency

---
 .../triton_kernel/shared_expert_gate.py       |  5 +++-
 .../qwen3next/test_shared_expert_gate.py      | 23 +++++++++++++++++++
 2 files changed, 27 insertions(+), 1 deletion(-)
 create mode 100644 unit_tests/models/qwen3next/test_shared_expert_gate.py

diff --git a/lightllm/models/qwen3next/triton_kernel/shared_expert_gate.py b/lightllm/models/qwen3next/triton_kernel/shared_expert_gate.py
index 8b73cfd74d..c9872a35ad 100644
--- a/lightllm/models/qwen3next/triton_kernel/shared_expert_gate.py
+++ b/lightllm/models/qwen3next/triton_kernel/shared_expert_gate.py
@@ -24,7 +24,10 @@ def _sigmoid_mul_kernel(
         gate_vals = tl.load(gate + row * stride_g_m).to(tl.float32)
     else:
         gate_vals = tl.load(gate + row * stride_g_m + offs * stride_g_n, mask=mask, other=0.0).to(tl.float32)
-    gate_vals = tl.sigmoid(gate_vals)
+    # Match the unfused PyTorch path used by training: sigmoid is materialized
+    # in the gate dtype before the multiplication. Keeping this intermediate
+    # rounding is important for rollout/training logprob consistency in BF16.
+    gate_vals = tl.sigmoid(gate_vals).to(gate.dtype.element_ty).to(tl.float32)
     tl.store(x_ptrs, (x_vals * gate_vals).to(x.dtype.element_ty), mask=mask)
 
 
diff --git a/unit_tests/models/qwen3next/test_shared_expert_gate.py b/unit_tests/models/qwen3next/test_shared_expert_gate.py
new file mode 100644
index 0000000000..e57c0083a7
--- /dev/null
+++ b/unit_tests/models/qwen3next/test_shared_expert_gate.py
@@ -0,0 +1,23 @@
+import pytest
+import torch
+
+from lightllm.models.qwen3next.triton_kernel.shared_expert_gate import sigmoid_mul_
+
+
+@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float16])
+@pytest.mark.parametrize("gate_width", [1, 6144])
+def test_sigmoid_mul_matches_unfused_dtype_rounding(dtype, gate_width):
+    if not torch.cuda.is_available():
+        pytest.skip("CUDA is required")
+
+    torch.manual_seed(2026)
+    x = torch.randn((17, 6144), device="cuda", dtype=dtype) * 2
+    gate = torch.randn((17, gate_width), device="cuda", dtype=dtype) * 3
+
+    expected = x.clone()
+    expected.mul_(gate.clone().sigmoid_())
+
+    actual = x.clone()
+    sigmoid_mul_(actual, gate)
+
+    torch.testing.assert_close(actual, expected, rtol=0, atol=0)

From 2798d369d9957d442f8cb40a1bc50535f75b0589 Mon Sep 17 00:00:00 2001
From: shihaobai <1798930569@qq.com>
Date: Thu, 30 Jul 2026 05:11:29 +0000
Subject: [PATCH 2/9] clarify qwen3.5 gate rounding reference

---
 .../models/qwen3next/triton_kernel/shared_expert_gate.py     | 5 ++---
 1 file changed, 2 insertions(+), 3 deletions(-)

diff --git a/lightllm/models/qwen3next/triton_kernel/shared_expert_gate.py b/lightllm/models/qwen3next/triton_kernel/shared_expert_gate.py
index c9872a35ad..b9c0c7b4af 100644
--- a/lightllm/models/qwen3next/triton_kernel/shared_expert_gate.py
+++ b/lightllm/models/qwen3next/triton_kernel/shared_expert_gate.py
@@ -24,9 +24,8 @@ def _sigmoid_mul_kernel(
         gate_vals = tl.load(gate + row * stride_g_m).to(tl.float32)
     else:
         gate_vals = tl.load(gate + row * stride_g_m + offs * stride_g_n, mask=mask, other=0.0).to(tl.float32)
-    # Match the unfused PyTorch path used by training: sigmoid is materialized
-    # in the gate dtype before the multiplication. Keeping this intermediate
-    # rounding is important for rollout/training logprob consistency in BF16.
+    # Match LightLLM's pre-fusion sigmoid_() then mul_() behavior: sigmoid is
+    # materialized in the gate dtype before the multiplication.
     gate_vals = tl.sigmoid(gate_vals).to(gate.dtype.element_ty).to(tl.float32)
     tl.store(x_ptrs, (x_vals * gate_vals).to(x.dtype.element_ty), mask=mask)
 

From 537a584d0915e616003809d05d126f4f26194880 Mon Sep 17 00:00:00 2001
From: shihaobai <1798930569@qq.com>
Date: Thu, 30 Jul 2026 07:43:12 +0000
Subject: [PATCH 3/9] fix qwen3.5 moe RL expert weight splitting

---
 .../layer_weights/transformer_layer_weight.py | 22 +++++++-
 .../models/qwen3_5_moe/test_weight_split.py   | 50 +++++++++++++++++++
 2 files changed, 70 insertions(+), 2 deletions(-)
 create mode 100644 unit_tests/models/qwen3_5_moe/test_weight_split.py

diff --git a/lightllm/models/qwen3_5_moe/layer_weights/transformer_layer_weight.py b/lightllm/models/qwen3_5_moe/layer_weights/transformer_layer_weight.py
index 7a91ed47e3..6ed23dae82 100644
--- a/lightllm/models/qwen3_5_moe/layer_weights/transformer_layer_weight.py
+++ b/lightllm/models/qwen3_5_moe/layer_weights/transformer_layer_weight.py
@@ -1,4 +1,6 @@
-from lightllm.models.qwen3_5.layer_weights.transformer_layer_weight import Qwen35TransformerLayerWeight
+from lightllm.models.qwen3_5.layer_weights.transformer_layer_weight import (
+    Qwen35TransformerLayerWeight,
+)
 
 
 class Qwen35MOETransformerLayerWeight(Qwen35TransformerLayerWeight):
@@ -11,7 +13,9 @@ def split_fused_expert_weights(weights: dict, layer_num: int, moe_intermediate_s
     """将 HF 打包的 fused MoE expert 权重拆成按 expert 索引的独立权重。
 
     部分 checkpoint（如 Qwen3.5-MoE）把所有 expert 的 gate_up / down 压成
-    ``mlp.experts.{gate_up,down}_proj`` 的打包张量（首维为 expert 数）。
+    ``mlp.experts.{gate_up,down}_proj`` 的打包张量（首维为 expert 数）；
+    在线 RL 更新也可能传入独立的 ``gate_proj`` / ``up_proj`` / ``down_proj``
+    三个 expert-packed 张量。
     本函数只处理 ``model.layers.{layer_num}`` 下的这类 key：弹出打包权重，
     再写入 ``mlp.experts.{expert_idx}.{gate,up,down}_proj.weight``，供后续
     按 expert 加载。``gate_up_proj`` 会按 ``moe_intermediate_size`` 沿
@@ -34,6 +38,20 @@ def split_fused_expert_weights(weights: dict, layer_num: int, moe_intermediate_s
                 weights[f"{prefix}.{expert_idx}.gate_proj.weight"] = gate_weight[expert_idx]
                 weights[f"{prefix}.{expert_idx}.up_proj.weight"] = up_weight[expert_idx]
 
+        elif "mlp.experts.gate_proj" in k:
+            gate_weight = weights.pop(k)
+            prefix = k.rsplit(".gate_proj", 1)[0]
+
+            for expert_idx in range(gate_weight.shape[0]):
+                weights[f"{prefix}.{expert_idx}.gate_proj.weight"] = gate_weight[expert_idx]
+
+        elif "mlp.experts.up_proj" in k:
+            up_weight = weights.pop(k)
+            prefix = k.rsplit(".up_proj", 1)[0]
+
+            for expert_idx in range(up_weight.shape[0]):
+                weights[f"{prefix}.{expert_idx}.up_proj.weight"] = up_weight[expert_idx]
+
         elif "mlp.experts.down_proj" in k:
             down_weight = weights.pop(k)
             prefix = k.rsplit(".down_proj", 1)[0]
diff --git a/unit_tests/models/qwen3_5_moe/test_weight_split.py b/unit_tests/models/qwen3_5_moe/test_weight_split.py
new file mode 100644
index 0000000000..74229c7029
--- /dev/null
+++ b/unit_tests/models/qwen3_5_moe/test_weight_split.py
@@ -0,0 +1,50 @@
+import torch
+
+from lightllm.models.qwen3_5_moe.layer_weights.transformer_layer_weight import (
+    split_fused_expert_weights,
+)
+
+
+def test_split_independent_expert_packed_weights_from_rl_update():
+    num_experts = 3
+    intermediate_size = 2
+    hidden_size = 4
+    prefix = "model.layers.7.mlp.experts"
+    gate = torch.arange(num_experts * intermediate_size * hidden_size).view(num_experts, intermediate_size, hidden_size)
+    up = gate + 100
+    down = torch.arange(num_experts * hidden_size * intermediate_size).view(num_experts, hidden_size, intermediate_size)
+    weights = {
+        f"{prefix}.gate_proj.weight": gate,
+        f"{prefix}.up_proj.weight": up,
+        f"{prefix}.down_proj.weight": down,
+        "model.layers.8.mlp.experts.gate_proj.weight": torch.empty(0),
+    }
+
+    split_fused_expert_weights(weights, layer_num=7, moe_intermediate_size=intermediate_size)
+
+    assert f"{prefix}.gate_proj.weight" not in weights
+    assert f"{prefix}.up_proj.weight" not in weights
+    assert f"{prefix}.down_proj.weight" not in weights
+    assert "model.layers.8.mlp.experts.gate_proj.weight" in weights
+    for expert_idx in range(num_experts):
+        torch.testing.assert_close(weights[f"{prefix}.{expert_idx}.gate_proj.weight"], gate[expert_idx])
+        torch.testing.assert_close(weights[f"{prefix}.{expert_idx}.up_proj.weight"], up[expert_idx])
+        torch.testing.assert_close(weights[f"{prefix}.{expert_idx}.down_proj.weight"], down[expert_idx])
+
+
+def test_split_checkpoint_gate_up_expert_packed_weight():
+    num_experts = 3
+    intermediate_size = 2
+    hidden_size = 4
+    prefix = "model.layers.7.mlp.experts"
+    gate = torch.arange(num_experts * intermediate_size * hidden_size).view(num_experts, intermediate_size, hidden_size)
+    up = gate + 100
+    gate_up = torch.cat((gate, up), dim=1)
+    weights = {f"{prefix}.gate_up_proj": gate_up}
+
+    split_fused_expert_weights(weights, layer_num=7, moe_intermediate_size=intermediate_size)
+
+    assert f"{prefix}.gate_up_proj" not in weights
+    for expert_idx in range(num_experts):
+        torch.testing.assert_close(weights[f"{prefix}.{expert_idx}.gate_proj.weight"], gate[expert_idx])
+        torch.testing.assert_close(weights[f"{prefix}.{expert_idx}.up_proj.weight"], up[expert_idx])

From 01143a32d0582d02c23df37f31e93817755bda36 Mon Sep 17 00:00:00 2001
From: shihaobai <1798930569@qq.com>
Date: Thu, 30 Jul 2026 08:53:33 +0000
Subject: [PATCH 4/9] simplify qwen3.5 BF16 gate multiply

---
 .../models/qwen3next/triton_kernel/shared_expert_gate.py | 9 +++++++--
 1 file changed, 7 insertions(+), 2 deletions(-)

diff --git a/lightllm/models/qwen3next/triton_kernel/shared_expert_gate.py b/lightllm/models/qwen3next/triton_kernel/shared_expert_gate.py
index b9c0c7b4af..991f06ddcf 100644
--- a/lightllm/models/qwen3next/triton_kernel/shared_expert_gate.py
+++ b/lightllm/models/qwen3next/triton_kernel/shared_expert_gate.py
@@ -19,14 +19,19 @@ def _sigmoid_mul_kernel(
     offs = tl.arange(0, BLOCK_N)
     mask = offs < N
     x_ptrs = x + row * stride_x_m + offs * stride_x_n
-    x_vals = tl.load(x_ptrs, mask=mask, other=0.0).to(tl.float32)
+    x_vals = tl.load(x_ptrs, mask=mask, other=0.0)
     if GATE_N == 1:
         gate_vals = tl.load(gate + row * stride_g_m).to(tl.float32)
     else:
         gate_vals = tl.load(gate + row * stride_g_m + offs * stride_g_n, mask=mask, other=0.0).to(tl.float32)
     # Match LightLLM's pre-fusion sigmoid_() then mul_() behavior: sigmoid is
     # materialized in the gate dtype before the multiplication.
-    gate_vals = tl.sigmoid(gate_vals).to(gate.dtype.element_ty).to(tl.float32)
+    gate_vals = tl.sigmoid(gate_vals).to(gate.dtype.element_ty)
+    # Qwen3.5 uses BF16, so keep both rounded operands in BF16 for the direct
+    # multiply. Preserve the widened multiply for the generic FP16 path.
+    if x.dtype.element_ty == tl.float16:
+        x_vals = x_vals.to(tl.float32)
+        gate_vals = gate_vals.to(tl.float32)
     tl.store(x_ptrs, (x_vals * gate_vals).to(x.dtype.element_ty), mask=mask)
 
 

From 49321c637cbf101d240496b1704a600c16a50cc0 Mon Sep 17 00:00:00 2001
From: shihaobai <1798930569@qq.com>
Date: Fri, 31 Jul 2026 07:36:30 +0000
Subject: [PATCH 5/9] fix(qwen3.5): restore split qkv and gate projections

---
 .../layer_infer/transformer_layer_infer.py    | 10 ++-------
 .../layer_infer/transformer_layer_infer.py    | 10 ++-------
 .../layer_weights/transformer_layer_weight.py | 21 ++++++++++++-------
 3 files changed, 18 insertions(+), 23 deletions(-)

diff --git a/lightllm/models/qwen3_5/layer_infer/transformer_layer_infer.py b/lightllm/models/qwen3_5/layer_infer/transformer_layer_infer.py
index 649db03b11..c9f4078998 100644
--- a/lightllm/models/qwen3_5/layer_infer/transformer_layer_infer.py
+++ b/lightllm/models/qwen3_5/layer_infer/transformer_layer_infer.py
@@ -28,17 +28,11 @@ def _get_qkv(
         input = input.view(-1, self.embed_dim_)
         input = self._tpsp_allgather(input=input, infer_state=infer_state)
 
-        qkv_gate_out = layer_weight.qkvo_gate_proj.mm(input)
-        qkv_out, o_gate = qkv_gate_out.split(
-            [
-                self.tp_q_head_num_ * self.head_dim_ + (self.tp_k_head_num_ + self.tp_v_head_num_) * self.head_dim_,
-                self.tp_q_head_num_ * self.head_dim_,
-            ],
-            dim=-1,
-        )
+        qkv_out = layer_weight.qkv_proj.mm(input)
         q, cache_kv = qkv_out.split(
             [self.tp_q_head_num_ * self.head_dim_, (self.tp_k_head_num_ + self.tp_v_head_num_) * self.head_dim_], dim=-1
         )
+        o_gate = layer_weight._o_gate_proj.mm(input)
 
         infer_state.gate_logics_value = o_gate
         layer_weight.qk_norm_weight_(
diff --git a/lightllm/models/qwen3next/layer_infer/transformer_layer_infer.py b/lightllm/models/qwen3next/layer_infer/transformer_layer_infer.py
index 92d68c9fd2..d388f701a9 100644
--- a/lightllm/models/qwen3next/layer_infer/transformer_layer_infer.py
+++ b/lightllm/models/qwen3next/layer_infer/transformer_layer_infer.py
@@ -131,18 +131,12 @@ def _get_qkv(
     ) -> Tuple[torch.Tensor, torch.Tensor]:
         input = input.view(-1, self.embed_dim_)
         input = self._tpsp_allgather(input=input, infer_state=infer_state)
-        qkv_gate_out = layer_weight.qkvo_gate_proj.mm(input)
-        qkv_out, o_gate = qkv_gate_out.split(
-            [
-                self.tp_q_head_num_ * self.head_dim_ * 2 + (self.tp_k_head_num_ + self.tp_v_head_num_) * self.head_dim_,
-                self.tp_q_head_num_ * self.head_dim_,
-            ],
-            dim=-1,
-        )
+        qkv_out = layer_weight.qkv_proj.mm(input)
         q, cache_kv = qkv_out.split(
             [self.tp_q_head_num_ * self.head_dim_ * 2, (self.tp_k_head_num_ + self.tp_v_head_num_) * self.head_dim_],
             dim=-1,
         )
+        o_gate = layer_weight._o_gate_proj.mm(input)
         infer_state.gate_logics_value = o_gate
         layer_weight.qk_norm_weight_(
             q,
diff --git a/lightllm/models/qwen3next/layer_weights/transformer_layer_weight.py b/lightllm/models/qwen3next/layer_weights/transformer_layer_weight.py
index 60901ad6b9..e3678da1cd 100644
--- a/lightllm/models/qwen3next/layer_weights/transformer_layer_weight.py
+++ b/lightllm/models/qwen3next/layer_weights/transformer_layer_weight.py
@@ -12,7 +12,6 @@
     QKGEMMANormWeight,
     FusedMoeWeight,
 )
-from lightllm.models.qwen3next.layer_weights.qkv_gated_rowmm_weight import QKVGatedROWNMMWeight
 
 
 class Qwen3NextTransformerLayerWeight(Qwen3MOETransformerLayerWeight):
@@ -24,17 +23,25 @@ def __init__(self, layer_num, data_type, network_config, quant_cfg=None):
 
     def _init_qkv(self):
         in_dim = self.n_embed
-        self._o_gate_weight_name = f"model.layers.{self.layer_num_}.self_attn.o_gate_proj.weight"
-        qkv_quant = self.get_quant_method("qkv_proj")
-        self.qkvo_gate_proj = QKVGatedROWNMMWeight(
+        q_out_dim = self.q_head_num_ * self.head_dim
+        self.qkv_proj = QKVROWNMMWeight(
             in_dim=in_dim,
             q_head_num=self.q_head_num_,
             kv_head_num=self.k_head_num_,
             head_dim=self.head_dim,
-            weight_names=[self._q_weight_name, self._k_weight_name, self._v_weight_name, self._o_gate_weight_name],
+            weight_names=[self._q_weight_name, self._k_weight_name, self._v_weight_name],
+            data_type=self.data_type_,
+            bias_names=[self._q_bias_name, self._k_bias_name, self._v_bias_name],
+            quant_method=self.get_quant_method("qkv_proj"),
+        )
+        self._o_gate_weight_name = f"model.layers.{self.layer_num_}.self_attn.o_gate_proj.weight"
+        self._o_gate_proj = ROWMMWeight(
+            in_dim=in_dim,
+            out_dims=[q_out_dim],
+            weight_names=[self._o_gate_weight_name],
             data_type=self.data_type_,
-            bias_names=[self._q_bias_name, self._k_bias_name, self._v_bias_name, None],
-            quant_method=qkv_quant,
+            bias_names=None,
+            quant_method=self.get_quant_method("o_gate_proj"),
         )
 
     def _init_weight(self):

From e89d9934c2454ea5630b2f27fa318d7d2eb44d44 Mon Sep 17 00:00:00 2001
From: shihaobai <1798930569@qq.com>
Date: Sun, 2 Aug 2026 09:59:29 +0000
Subject: [PATCH 6/9] fix: make gated RMSNorm launch deterministic

---
 .../triton_kernel/norm/gated_rmsnorm.py       |  65 ++--------
 ...torch.bfloat16}_NVIDIA_H100_80GB_HBM3.json | 118 ------------------
 ...6,x_dtype=torch.bfloat16}_NVIDIA_H200.json | 118 ------------------
 ...6,x_dtype=torch.bfloat16}_NVIDIA_H200.json |  50 --------
 4 files changed, 8 insertions(+), 343 deletions(-)
 delete mode 100644 lightllm/common/triton_utils/autotune_kernel_configs/triton_3.5.1/NVIDIA_H100_80GB_HBM3/gated_rmsnorm_forward:v1/{N=128,has_bias=false,weight_dtype=torch.bfloat16,x_dtype=torch.bfloat16}_NVIDIA_H100_80GB_HBM3.json
 delete mode 100644 lightllm/common/triton_utils/autotune_kernel_configs/triton_3.5.1/NVIDIA_H200/gated_rmsnorm_forward:v1/{N=128,has_bias=false,weight_dtype=torch.bfloat16,x_dtype=torch.bfloat16}_NVIDIA_H200.json
 delete mode 100644 lightllm/common/triton_utils/autotune_kernel_configs/triton_3.6.0/NVIDIA_H200/gated_rmsnorm_forward:v1/{N=128,has_bias=false,weight_dtype=torch.bfloat16,x_dtype=torch.bfloat16}_NVIDIA_H200.json

diff --git a/lightllm/common/basemodel/triton_kernel/norm/gated_rmsnorm.py b/lightllm/common/basemodel/triton_kernel/norm/gated_rmsnorm.py
index c62c5eb5d2..d2fca0a4c3 100644
--- a/lightllm/common/basemodel/triton_kernel/norm/gated_rmsnorm.py
+++ b/lightllm/common/basemodel/triton_kernel/norm/gated_rmsnorm.py
@@ -1,7 +1,6 @@
 import triton
 import triton.language as tl
 import torch
-from lightllm.common.triton_utils.autotuner import autotune
 
 
 @triton.heuristics(
@@ -60,39 +59,6 @@ def gated_rmsnorm_forward_kernel(
     tl.store(Y + cols, y, mask=mask)
 
 
-def _get_gated_rmsnorm_configs():
-    """Generate configurations for autotuning gated RMSNorm kernel."""
-    configs = []
-    # Different BLOCK_N sizes (powers of 2)
-    for block_n in [64, 128, 256, 512, 1024, 2048, 4096]:
-        # Different number of warps
-        for num_warps in [1, 2, 4, 8]:
-            # Skip configurations that are likely to be inefficient
-            if block_n >= 2048 and num_warps > 4:
-                continue
-            if block_n <= 128 and num_warps > 2:
-                continue
-            configs.append({"BLOCK_N": block_n, "num_warps": num_warps})
-    return configs
-
-
-def _get_gated_rmsnorm_static_key(x: torch.Tensor, weight: torch.Tensor, bias: torch.Tensor):
-    """Generate static key for caching autotuned configurations."""
-    M, N = x.shape
-    return {
-        "x_dtype": str(x.dtype),
-        "weight_dtype": str(weight.dtype),
-        "N": N,
-        "has_bias": bias is not None,
-    }
-
-
-@autotune(
-    kernel_name="gated_rmsnorm_forward:v1",
-    configs_gen_func=_get_gated_rmsnorm_configs,
-    static_key_func=_get_gated_rmsnorm_static_key,
-    run_key_func=lambda x: x.shape[0],
-)
 def gated_rmsnorm_forward(
     x: torch.Tensor,
     weight: torch.Tensor,
@@ -102,7 +68,6 @@ def gated_rmsnorm_forward(
     out: torch.Tensor = None,
     group_size: int = None,
     norm_before_gate: bool = True,
-    run_config: dict = None,
 ):
     M, N = x.shape
     if group_size is None:
@@ -125,27 +90,13 @@ def gated_rmsnorm_forward(
     else:
         out = torch.empty_like(x)
     assert out.stride(-1) == 1
-    # Default heuristic when autotune is disabled or no config provided
-    if not run_config:
-        # Less than 64KB per feature: enqueue fused kernel
-        MAX_FUSED_SIZE = 65536 // x.element_size()
-        BLOCK_N = min(MAX_FUSED_SIZE, triton.next_power_of_2(group_size))
-        if group_size > BLOCK_N:
-            raise RuntimeError("This layer norm doesn't support feature dim >= 64KB.")
-        # heuristics for number of warps
-        num_warps = min(max(BLOCK_N // 256, 1), 8)
-        run_config = {"BLOCK_N": BLOCK_N, "num_warps": num_warps}
-
-    BLOCK_N = run_config["BLOCK_N"]
-    num_warps = run_config["num_warps"]
-
-    # Validate BLOCK_N against group_size
-    if group_size > BLOCK_N:
-        # Fall back to largest valid BLOCK_N
-        MAX_FUSED_SIZE = 65536 // x.element_size()
-        BLOCK_N = min(MAX_FUSED_SIZE, triton.next_power_of_2(group_size))
-        if group_size > BLOCK_N:
-            raise RuntimeError("This layer norm doesn't support feature dim >= 64KB.")
+    # Keep the launch configuration deterministic: changing the reduction
+    # layout can shift token logprobs in recurrent Qwen3.5 GDN layers.
+    max_fused_size = 65536 // x.element_size()
+    block_n = min(max_fused_size, triton.next_power_of_2(group_size))
+    if group_size > block_n:
+        raise RuntimeError("This layer norm doesn't support feature dim >= 64KB.")
+    num_warps = min(max(block_n // 256, 1), 8)
 
     grid = (M, ngroups)
     gated_rmsnorm_forward_kernel[grid](
@@ -160,7 +111,7 @@ def gated_rmsnorm_forward(
         M,
         group_size,
         eps,
-        BLOCK_N=BLOCK_N,
+        BLOCK_N=block_n,
         NORM_BEFORE_GATE=norm_before_gate,
         num_warps=num_warps,
     )
diff --git a/lightllm/common/triton_utils/autotune_kernel_configs/triton_3.5.1/NVIDIA_H100_80GB_HBM3/gated_rmsnorm_forward:v1/{N=128,has_bias=false,weight_dtype=torch.bfloat16,x_dtype=torch.bfloat16}_NVIDIA_H100_80GB_HBM3.json b/lightllm/common/triton_utils/autotune_kernel_configs/triton_3.5.1/NVIDIA_H100_80GB_HBM3/gated_rmsnorm_forward:v1/{N=128,has_bias=false,weight_dtype=torch.bfloat16,x_dtype=torch.bfloat16}_NVIDIA_H100_80GB_HBM3.json
deleted file mode 100644
index 31d7a6e203..0000000000
--- a/lightllm/common/triton_utils/autotune_kernel_configs/triton_3.5.1/NVIDIA_H100_80GB_HBM3/gated_rmsnorm_forward:v1/{N=128,has_bias=false,weight_dtype=torch.bfloat16,x_dtype=torch.bfloat16}_NVIDIA_H100_80GB_HBM3.json
+++ /dev/null
@@ -1,118 +0,0 @@
-{
-  "1024": {
-    "BLOCK_N": 64,
-    "num_warps": 2
-  },
-  "12": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "1200": {
-    "BLOCK_N": 256,
-    "num_warps": 2
-  },
-  "12288": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "128": {
-    "BLOCK_N": 256,
-    "num_warps": 8
-  },
-  "131072": {
-    "BLOCK_N": 64,
-    "num_warps": 2
-  },
-  "1536": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "16": {
-    "BLOCK_N": 64,
-    "num_warps": 2
-  },
-  "1600": {
-    "BLOCK_N": 64,
-    "num_warps": 1
-  },
-  "16384": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "192": {
-    "BLOCK_N": 512,
-    "num_warps": 1
-  },
-  "196608": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "2048": {
-    "BLOCK_N": 64,
-    "num_warps": 1
-  },
-  "24576": {
-    "BLOCK_N": 64,
-    "num_warps": 1
-  },
-  "256": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "262144": {
-    "BLOCK_N": 64,
-    "num_warps": 2
-  },
-  "3072": {
-    "BLOCK_N": 256,
-    "num_warps": 2
-  },
-  "32768": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "384": {
-    "BLOCK_N": 512,
-    "num_warps": 2
-  },
-  "4096": {
-    "BLOCK_N": 256,
-    "num_warps": 1
-  },
-  "49152": {
-    "BLOCK_N": 64,
-    "num_warps": 2
-  },
-  "512": {
-    "BLOCK_N": 64,
-    "num_warps": 2
-  },
-  "64": {
-    "BLOCK_N": 256,
-    "num_warps": 8
-  },
-  "65536": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "768": {
-    "BLOCK_N": 256,
-    "num_warps": 2
-  },
-  "8": {
-    "BLOCK_N": 64,
-    "num_warps": 2
-  },
-  "800": {
-    "BLOCK_N": 64,
-    "num_warps": 2
-  },
-  "8192": {
-    "BLOCK_N": 256,
-    "num_warps": 1
-  },
-  "96": {
-    "BLOCK_N": 512,
-    "num_warps": 2
-  }
-}
\ No newline at end of file
diff --git a/lightllm/common/triton_utils/autotune_kernel_configs/triton_3.5.1/NVIDIA_H200/gated_rmsnorm_forward:v1/{N=128,has_bias=false,weight_dtype=torch.bfloat16,x_dtype=torch.bfloat16}_NVIDIA_H200.json b/lightllm/common/triton_utils/autotune_kernel_configs/triton_3.5.1/NVIDIA_H200/gated_rmsnorm_forward:v1/{N=128,has_bias=false,weight_dtype=torch.bfloat16,x_dtype=torch.bfloat16}_NVIDIA_H200.json
deleted file mode 100644
index 55ccb24a65..0000000000
--- a/lightllm/common/triton_utils/autotune_kernel_configs/triton_3.5.1/NVIDIA_H200/gated_rmsnorm_forward:v1/{N=128,has_bias=false,weight_dtype=torch.bfloat16,x_dtype=torch.bfloat16}_NVIDIA_H200.json
+++ /dev/null
@@ -1,118 +0,0 @@
-{
-  "1024": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "128": {
-    "BLOCK_N": 512,
-    "num_warps": 2
-  },
-  "131072": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "1536": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "16": {
-    "BLOCK_N": 256,
-    "num_warps": 4
-  },
-  "1600": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "16384": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "192": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "2048": {
-    "BLOCK_N": 64,
-    "num_warps": 2
-  },
-  "24": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "2400": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "24576": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "256": {
-    "BLOCK_N": 512,
-    "num_warps": 2
-  },
-  "262144": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "3072": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "32768": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "384": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "393216": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "4096": {
-    "BLOCK_N": 64,
-    "num_warps": 2
-  },
-  "49152": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "512": {
-    "BLOCK_N": 256,
-    "num_warps": 4
-  },
-  "6144": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "64": {
-    "BLOCK_N": 64,
-    "num_warps": 2
-  },
-  "65536": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  },
-  "768": {
-    "BLOCK_N": 256,
-    "num_warps": 2
-  },
-  "8": {
-    "BLOCK_N": 64,
-    "num_warps": 2
-  },
-  "800": {
-    "BLOCK_N": 64,
-    "num_warps": 2
-  },
-  "8192": {
-    "BLOCK_N": 128,
-    "num_warps": 2
-  },
-  "98304": {
-    "BLOCK_N": 128,
-    "num_warps": 1
-  }
-}
\ No newline at end of file
diff --git a/lightllm/common/triton_utils/autotune_kernel_configs/triton_3.6.0/NVIDIA_H200/gated_rmsnorm_forward:v1/{N=128,has_bias=false,weight_dtype=torch.bfloat16,x_dtype=torch.bfloat16}_NVIDIA_H200.json b/lightllm/common/triton_utils/autotune_kernel_configs/triton_3.6.0/NVIDIA_H200/gated_rmsnorm_forward:v1/{N=128,has_bias=false,weight_dtype=torch.bfloat16,x_dtype=torch.bfloat16}_NVIDIA_H200.json
deleted file mode 100644
index 1fcfa30e97..0000000000
--- a/lightllm/common/triton_utils/autotune_kernel_configs/triton_3.6.0/NVIDIA_H200/gated_rmsnorm_forward:v1/{N=128,has_bias=false,weight_dtype=torch.bfloat16,x_dtype=torch.bfloat16}_NVIDIA_H200.json
+++ /dev/null
@@ -1,50 +0,0 @@
-{
-  "1024": {
-    "BLOCK_N": 128,
-    "num_warps": 2
-  },
-  "128": {
-    "BLOCK_N": 128,
-    "num_warps": 2
-  },
-  "131072": {
-    "BLOCK_N": 128,
-    "num_warps": 2
-  },
-  "16384": {
-    "BLOCK_N": 128,
-    "num_warps": 2
-  },
-  "2048": {
-    "BLOCK_N": 128,
-    "num_warps": 2
-  },
-  "256": {
-    "BLOCK_N": 128,
-    "num_warps": 2
-  },
-  "32768": {
-    "BLOCK_N": 128,
-    "num_warps": 2
-  },
-  "512": {
-    "BLOCK_N": 128,
-    "num_warps": 2
-  },
-  "64": {
-    "BLOCK_N": 512,
-    "num_warps": 4
-  },
-  "8": {
-    "BLOCK_N": 128,
-    "num_warps": 2
-  },
-  "800": {
-    "BLOCK_N": 128,
-    "num_warps": 2
-  },
-  "8192": {
-    "BLOCK_N": 128,
-    "num_warps": 2
-  }
-}
\ No newline at end of file

From 7a38ac12c8cf9de6fe1ddcd789323bb61b3b63a5 Mon Sep 17 00:00:00 2001
From: shihaobai <1798930569@qq.com>
Date: Sun, 2 Aug 2026 13:27:09 +0000
Subject: [PATCH 7/9] fix: restore multimodal image pixel default

---
 lightllm/server/api_cli.py | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/lightllm/server/api_cli.py b/lightllm/server/api_cli.py
index aea6eb7305..0c96630e51 100644
--- a/lightllm/server/api_cli.py
+++ b/lightllm/server/api_cli.py
@@ -447,7 +447,7 @@ def add_cli_args(parser: argparse.ArgumentParser) -> argparse.ArgumentParser:
     parser.add_argument(
         "--max_image_pixels",
         type=int,
-        default=3686400,  # 8294400 is 4k, 3686400 is 2k
+        default=8294400,
         help="maximum allowed pixel count for one image before resize preprocessing",
     )
     parser.add_argument(

From 218214feb6090115e49157389ca69bef7657e9b0 Mon Sep 17 00:00:00 2001
From: shihaobai <1798930569@qq.com>
Date: Sun, 2 Aug 2026 14:05:45 +0000
Subject: [PATCH 8/9] fix: force NCCL all-reduce in RL mode

---
 lightllm/utils/dist_check_utils.py | 9 +++++++++
 1 file changed, 9 insertions(+)

diff --git a/lightllm/utils/dist_check_utils.py b/lightllm/utils/dist_check_utils.py
index 12b0b81993..c660b9bc4e 100644
--- a/lightllm/utils/dist_check_utils.py
+++ b/lightllm/utils/dist_check_utils.py
@@ -157,6 +157,15 @@ def auto_configure_allreduce_flags_from_args(args: "StartArgs") -> None:
 
     会就地修改 ``args.disable_flashinfer_allreduce`` / ``args.disable_symm_mem_allreduce``。
     """
+    if args.enable_rl:
+        # Keep rollout replicas on one deterministic communication path.  A
+        # capability probe can fail independently on different replicas and
+        # otherwise leave an RL job mixing custom all-reduce with NCCL.
+        logger.info("RL mode: force TP all-reduce to NCCL.")
+        args.disable_flashinfer_allreduce = True
+        args.disable_symm_mem_allreduce = True
+        return
+
     if not _should_run_allreduce_capability_check(args):
         return
 

From a881c16d24336ddfd4167e9a5731cedaecdf1bfc Mon Sep 17 00:00:00 2001
From: shihaobai <1798930569@qq.com>
Date: Sun, 2 Aug 2026 14:21:40 +0000
Subject: [PATCH 9/9] fix: align RL all-reduce policy

---
 lightllm/utils/dist_check_utils.py | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)

diff --git a/lightllm/utils/dist_check_utils.py b/lightllm/utils/dist_check_utils.py
index c660b9bc4e..b26889c7c7 100644
--- a/lightllm/utils/dist_check_utils.py
+++ b/lightllm/utils/dist_check_utils.py
@@ -158,12 +158,12 @@ def auto_configure_allreduce_flags_from_args(args: "StartArgs") -> None:
     会就地修改 ``args.disable_flashinfer_allreduce`` / ``args.disable_symm_mem_allreduce``。
     """
     if args.enable_rl:
-        # Keep rollout replicas on one deterministic communication path.  A
-        # capability probe can fail independently on different replicas and
-        # otherwise leave an RL job mixing custom all-reduce with NCCL.
-        logger.info("RL mode: force TP all-reduce to NCCL.")
+        # Match the established RL rollout policy on every replica.  Skipping
+        # capability probes also avoids per-replica divergence when concurrent
+        # local probes race for the same TCP port.
+        logger.info("RL mode: disable FlashInfer all-reduce and keep SymmMem all-reduce enabled.")
         args.disable_flashinfer_allreduce = True
-        args.disable_symm_mem_allreduce = True
+        args.disable_symm_mem_allreduce = False
         return
 
     if not _should_run_allreduce_capability_check(args):