fix(CRITICAL): align patch_ops.sh with comp 168 — keep base qwen3_5.py + upstream搬运

patch_ops.sh v2: conditional model layer deployment
搬运: moe_combine.cu, moe_compute_index.cu, fused_moe_xllm.cpp,
      qwen3_gated_delta_net_base.cpp/.h, ilu_layer_fused_moe.h, ilu_layer_attention.h

ex_engine/csrc/ilu/ixformer.h

@@ -1,4 +1,4 @@
-/* Copyright 2025 The xLLM Authors. All Rights Reserved.
+/* Copyright 2025-2026 The xLLM Authors.
 
 Licensed under the Apache License, Version 2.0 (the "License");
 you may not use this file except in compliance with the License.

ex_engine/csrc/ilu/utils.h

@@ -1,4 +1,4 @@
-/* Copyright 2025 The xLLM Authors. All Rights Reserved.
+/* Copyright 2025-2026 The xLLM Authors.
 
 Licensed under the Apache License, Version 2.0 (the "License");
 you may not use this file except in compliance with the License.

ex_engine/csrc/ilu_layer_attention.h

@@ -1,4 +1,4 @@
-/* Copyright 2025 The xLLM Authors. All Rights Reserved.
+/* Copyright 2025-2026 The xLLM Authors.
 
 Licensed under the Apache License, Version 2.0 (the "License");
 you may not use this file except in compliance with the License.

ex_engine/csrc/ilu_layer_fused_moe.h

@@ -1,4 +1,4 @@
-/* Copyright 2025 The xLLM Authors. All Rights Reserved.
+/* Copyright 2025-2026 The xLLM Authors.
 
 Licensed under the Apache License, Version 2.0 (the "License");
 you may not use this file except in compliance with the License.

ex_engine/csrc/moe/fused_moe_xllm.cpp

@@ -0,0 +1,124 @@
+/* Copyright 2025-2026 The xLLM Authors.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    https://github.com/jd-opensource/xllm/blob/main/LICENSE
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+==============================================================================*/
+
+#include "kernels/cuda/cuda_ops_api.h"
+#include "kernels/cuda/utils.h"
+#include "platform/device.h"
+#include "platform/platform.h"
+
+namespace xllm::kernel::cuda {
+
+torch::Tensor cutlass_fused_moe(
+    const torch::Tensor& input,                   // [num_tokens, hidden]
+    const torch::Tensor& token_selected_experts,  // [num_tokens, top_k]
+    const torch::Tensor& token_final_scales,      // [num_tokens, top_k]
+    const torch::Tensor&
+        fc1_expert_weights,  // [num_experts, inter_dim, hidden]
+    const torch::Tensor&
+        fc2_expert_weights,  // [num_experts, hidden, inter_dim]
+    torch::ScalarType output_dtype,
+    const std::vector<torch::Tensor>& quant_scales,
+    int32_t tp_size,
+    int32_t tp_rank,
+    int32_t ep_size,
+    int32_t ep_rank,
+    int32_t cluster_size,
+    int32_t cluster_rank,
+    const std::optional<torch::Tensor>& fc1_expert_biases,
+    const std::optional<torch::Tensor>& fc2_expert_biases,
+    const std::optional<torch::Tensor>& input_sf,
+    const std::optional<torch::Tensor>& swiglu_alpha,
+    const std::optional<torch::Tensor>& swiglu_beta,
+    const std::optional<torch::Tensor>& swiglu_limit,
+    const std::optional<torch::Tensor>& output,
+    bool enable_alltoall,
+    bool use_deepseek_fp8_block_scale,
+    bool use_w4_group_scaling,
+    bool use_mxfp8_act_scaling,
+    bool min_latency_mode,
+    bool use_packed_weights,
+    int32_t tune_max_num_tokens,
+    ActivationType activation_type) {
+  int64_t num_rows = input.size(0);
+  int64_t hidden_size = fc2_expert_weights.size(1);
+
+  if (min_latency_mode) {
+    num_rows *= fc2_expert_weights.size(0);
+  }
+
+  std::vector<int64_t> output_shape = {num_rows, hidden_size};
+  torch::Tensor result_output;
+  if (output.has_value() && output.value().defined()) {
+    result_output = output.value();
+  } else {
+    torch::TensorOptions options = input.options().dtype(output_dtype);
+    result_output = torch::empty(output_shape, options);
+  }
+
+  std::string fused_moe_uri = "fused_moe";
+  if (Platform::is_support_sm90a()) {
+    fused_moe_uri += "_90";
+  } else if (Platform::is_support_sm100a() || Platform::is_support_sm100f()) {
+    fused_moe_uri += "_100";
+  } else if (Platform::is_support_sm120a()) {
+    fused_moe_uri += "_120";
+  } else {
+    LOG(FATAL) << "FusedMoE is only supported on sm90, sm100, sm120.";
+  }
+
+  bind_tvmffi_stream_to_current_torch_stream(input.device());
+
+  ffi::Module fused_moe_runner =
+      get_function(fused_moe_uri, "init")(
+          to_dl_data_type(input.scalar_type()),
+          to_dl_data_type(fc1_expert_weights.scalar_type()),
+          to_dl_data_type(output_dtype),
+          use_deepseek_fp8_block_scale,
+          use_w4_group_scaling,
+          use_mxfp8_act_scaling,
+          use_packed_weights)
+          .cast<ffi::Module>();
+
+  fused_moe_runner->GetFunction("run_moe").value()(
+      to_ffi_tensor(result_output),
+      to_ffi_tensor(input),
+      to_ffi_tensor(token_selected_experts),
+      to_ffi_optional_tensor(token_final_scales),
+      to_ffi_tensor(fc1_expert_weights),
+      to_ffi_optional_tensor(fc1_expert_biases),
+      to_ffi_tensor(fc2_expert_weights),
+      to_ffi_optional_tensor(fc2_expert_biases),
+      to_ffi_optional_array_tensors(quant_scales),
+      to_ffi_optional_tensor(input_sf),
+      to_ffi_optional_tensor(swiglu_alpha),
+      to_ffi_optional_tensor(swiglu_beta),
+      to_ffi_optional_tensor(swiglu_limit),
+      tp_size,
+      tp_rank,
+      ep_size,
+      ep_rank,
+      cluster_size,
+      cluster_rank,
+      enable_alltoall,
+      min_latency_mode,
+      /*profile_ids=*/ffi::Optional<ffi::Array<int64_t>>(),  // TODO: support
+                                                             // auto tuning
+                                                             // profile ids
+      support_pdl(),
+      activation_type);
+
+  return result_output;
+}
+}  // namespace xllm::kernel::cuda

ex_engine/csrc/moe/moe_combine.cu

@@ -0,0 +1,105 @@
+/* Copyright 2025-2026 The xLLM Authors. All Rights Reserved.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    https://github.com/jd-opensource/xllm/blob/main/LICENSE
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+==============================================================================*/
+
+// Fused MoE combine kernel — reorder + weighted sum in one pass.
+// Replaces: torch::zeros + index_copy_ + view + multiply + sum
+//
+// Algorithm per token (each block handles one token):
+//   1. For each of its topk experts, read gemm2 at flat_idx directly
+//      (gemm2 is flat-index-ordered after scatter via index_copy_ with dst_src)
+//   2. Multiply by router weight
+//   3. Accumulate into output[token]
+//
+// Grid:  num_tokens (N) blocks
+// Block: HIDDEN_DIM / HIDDEN_TILE threads
+
+#include <c10/cuda/CUDAGuard.h>
+
+#include "device_utils.cuh"
+#include "kernels/cuda/cuda_ops_api.h"
+
+namespace xllm::kernel::cuda {
+
+constexpr int32_t kCombineBlockSize = 256;
+
+template <typename scalar_t>
+__global__ void XLLM_KERNEL_ATTR(kCombineBlockSize) moe_combine_kernel(
+    const scalar_t* __restrict__ gemm2,       // [N*topk, H] flat-index-ordered
+    const float* __restrict__ reduce_weight,  // [N, topk]
+    scalar_t* __restrict__ output,            // [N, H]
+    int64_t N,
+    int32_t topk,
+    int64_t H) {
+  int64_t token_id = blockIdx.x;  // 0 .. N-1
+  if (token_id >= N) return;
+
+  int32_t tid = threadIdx.x;
+  int32_t stride = kCombineBlockSize;
+
+  // Accumulate over topk experts for this token
+  for (int64_t h = tid; h < H; h += stride) {
+    float acc = 0.0f;
+    for (int32_t k = 0; k < topk; ++k) {
+      int64_t flat_idx = token_id * topk + k;
+      float w = reduce_weight[flat_idx];
+      acc += w * static_cast<float>(gemm2[flat_idx * H + h]);
+    }
+    output[token_id * H + h] = static_cast<scalar_t>(acc);
+  }
+}
+
+// ---- Host-side orchestrator ----
+torch::Tensor moe_combine_result(
+    const torch::Tensor& gemm2,          // [N*topk, H] flat-index-ordered
+    const torch::Tensor& reduce_weight,  // [N, topk] float or same as gemm2
+    int64_t N,
+    int32_t topk) {
+  auto stream = at::cuda::getCurrentCUDAStream();
+  int64_t H = gemm2.size(1);
+  auto dtype = gemm2.scalar_type();
+
+  auto output = torch::empty({N, H}, gemm2.options());
+  auto rw = reduce_weight.to(gemm2.device(), torch::kFloat32).contiguous();
+
+  if (dtype == torch::kFloat16) {
+    moe_combine_kernel<c10::Half>
+        <<<N, kCombineBlockSize, 0, stream>>>(gemm2.data_ptr<c10::Half>(),
+                                              rw.data_ptr<float>(),
+                                              output.data_ptr<c10::Half>(),
+                                              N,
+                                              topk,
+                                              H);
+  } else if (dtype == torch::kBFloat16) {
+    moe_combine_kernel<c10::BFloat16>
+        <<<N, kCombineBlockSize, 0, stream>>>(gemm2.data_ptr<c10::BFloat16>(),
+                                              rw.data_ptr<float>(),
+                                              output.data_ptr<c10::BFloat16>(),
+                                              N,
+                                              topk,
+                                              H);
+  } else {
+    moe_combine_kernel<float>
+        <<<N, kCombineBlockSize, 0, stream>>>(gemm2.data_ptr<float>(),
+                                              rw.data_ptr<float>(),
+                                              output.data_ptr<float>(),
+                                              N,
+                                              topk,
+                                              H);
+  }
+
+  return output;
+}
+
+}  // namespace xllm::kernel::cuda

ex_engine/csrc/moe/moe_compute_index.cu

@@ -0,0 +1,155 @@
+/* Copyright 2025-2026 The xLLM Authors. All Rights Reserved.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    https://github.com/jd-opensource/xllm/blob/main/LICENSE
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+==============================================================================*/
+
+// Fused MoE token index computation — 3 kernels replacing:
+//   torch::bincount + 2 × torch::argsort + torch::cumsum + CPU sync
+//
+// Phase 1  histogram:      atomicAdd per-expert token counts
+// Phase 2  prefix_sum:     1 block, exclusive scan → expert_offsets
+// Phase 3  place_indices:  atomicAdd on offsets, write dst_src + src_dst
+//
+// expert_sizes  = per-expert token count [num_experts]  (preserved)
+// expert_offsets = exclusive prefix sum of counts       (scratch, reused)
+
+#include <c10/cuda/CUDAGuard.h>
+
+#include <cub/block/block_scan.cuh>
+
+#include "kernels/cuda/cuda_ops_api.h"
+
+namespace xllm::kernel::cuda {
+
+constexpr int32_t kMoeIndexBlock = 256;
+
+// ---- Phase 1: histogram ----
+__global__ void
+#ifdef USE_DCU
+__launch_bounds__(kMoeIndexBlock, 1)
+#endif
+    moe_histogram_kernel(const int32_t* __restrict__ expert_id,
+                         int32_t* __restrict__ expert_sizes,
+                         int64_t num_elements,
+                         int32_t num_experts) {
+  int64_t tid = int64_t(blockIdx.x) * kMoeIndexBlock + threadIdx.x;
+  if (tid < num_elements) {
+    int32_t eid = expert_id[tid];
+    if (eid >= 0 && eid < num_experts) {
+      atomicAdd(&expert_sizes[eid], 1);
+    }
+  }
+}
+
+// ---- Phase 2: exclusive prefix sum (1 block) ----
+// input:  expert_sizes  (per-expert counts)
+// output: expert_offsets (exclusive scan of counts)
+//         total_out      (total number of tokens, scalar)
+__global__ void
+#ifdef USE_DCU
+__launch_bounds__(kMoeIndexBlock, 1)
+#endif
+    moe_prefix_sum_kernel(const int32_t* __restrict__ expert_sizes,
+                          int32_t* __restrict__ expert_offsets,
+                          int32_t num_experts,
+                          int64_t* __restrict__ total_out) {
+  using BlockScan = cub::BlockScan<int32_t, kMoeIndexBlock>;
+  __shared__ typename BlockScan::TempStorage s_scan;
+
+  int32_t val = (threadIdx.x < num_experts) ? expert_sizes[threadIdx.x] : 0;
+  int32_t offset;
+  BlockScan(s_scan).ExclusiveSum(val, offset);
+  __syncthreads();
+
+  // total = all elements sum = last thread's exclusive output + its input
+  int32_t total = offset + val;
+
+  if (threadIdx.x < num_experts) {
+    expert_offsets[threadIdx.x] = offset;
+  }
+  if (threadIdx.x == 0 && total_out != nullptr) {
+    *total_out = total;
+  }
+}
+
+// ---- Phase 3: place indices ----
+// atomicAdd on expert_offsets to assign a unique position within
+// [start(e), start(e)+count(e)), then write both direction mappings.
+__global__ void
+#ifdef USE_DCU
+__launch_bounds__(kMoeIndexBlock, 1)
+#endif
+    moe_place_indices_kernel(const int32_t* __restrict__ expert_id,
+                             int32_t* __restrict__ expert_offsets,
+                             int32_t* __restrict__ dst_src,
+                             int32_t* __restrict__ src_dst,
+                             int64_t num_elements,
+                             int32_t num_experts) {
+  int64_t flat_idx = int64_t(blockIdx.x) * kMoeIndexBlock + threadIdx.x;
+  if (flat_idx >= num_elements) return;
+
+  int32_t eid = expert_id[flat_idx];
+  if (eid < 0 || eid >= num_experts) return;
+
+  int32_t pos = atomicAdd(&expert_offsets[eid], 1);
+  dst_src[pos] = static_cast<int32_t>(flat_idx);
+  src_dst[flat_idx] = pos;
+}
+
+// ---- Host-side orchestrator ----
+// Returns {src_dst, dst_src, expert_sizes}
+std::tuple<torch::Tensor, torch::Tensor, torch::Tensor> moe_compute_index(
+    const torch::Tensor& expert_id,
+    int64_t num_experts) {
+  auto device = expert_id.device();
+  auto stream = at::cuda::getCurrentCUDAStream();
+  int64_t N = expert_id.numel();
+  int32_t E = static_cast<int32_t>(num_experts);
+  CHECK_LE(E, kMoeIndexBlock) << "num_experts cannot exceed " << kMoeIndexBlock;
+  auto expert_id_i32 = expert_id.to(torch::kInt32).contiguous();
+  auto opt_i32 = expert_id_i32.options();
+
+  auto expert_sizes = torch::zeros({num_experts}, opt_i32);
+  auto expert_offsets = torch::empty({num_experts}, opt_i32);
+  auto dst_src = torch::empty({N}, opt_i32);
+  auto src_dst = torch::empty({N}, opt_i32);
+
+  int64_t grid = (N + kMoeIndexBlock - 1) / kMoeIndexBlock;
+
+  // Phase 1: histogram
+  moe_histogram_kernel<<<grid, kMoeIndexBlock, 0, stream>>>(
+      expert_id_i32.data_ptr<int32_t>(),
+      expert_sizes.data_ptr<int32_t>(),
+      N,
+      E);
+
+  // Phase 2: prefix sum (1 block)
+  moe_prefix_sum_kernel<<<1, kMoeIndexBlock, 0, stream>>>(
+      expert_sizes.data_ptr<int32_t>(),
+      expert_offsets.data_ptr<int32_t>(),
+      E,
+      nullptr);
+
+  // Phase 3: place indices
+  moe_place_indices_kernel<<<grid, kMoeIndexBlock, 0, stream>>>(
+      expert_id_i32.data_ptr<int32_t>(),
+      expert_offsets.data_ptr<int32_t>(),
+      dst_src.data_ptr<int32_t>(),
+      src_dst.data_ptr<int32_t>(),
+      N,
+      E);
+
+  return std::make_tuple(src_dst, dst_src, expert_sizes);
+}
+
+}  // namespace xllm::kernel::cuda

ex_engine/csrc/qwen3_gated_delta_net_base.h

@@ -0,0 +1,112 @@
+/* Copyright 2025-2026 The xLLM Authors.
+
+Licensed under the Apache License, Version 2.0 (the "License");
+you may not use this file except in compliance with the License.
+You may obtain a copy of the License at
+
+    https://github.com/jd-opensource/xllm/blob/main/LICENSE
+
+Unless required by applicable law or agreed to in writing, software
+distributed under the License is distributed on an "AS IS" BASIS,
+WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+See the License for the specific language governing permissions and
+limitations under the License.
+==============================================================================*/
+
+#pragma once
+
+#include <torch/torch.h>
+
+#include <optional>
+#include <string>
+#include <tuple>
+#include <utility>
+
+#include "attention.h"
+#include "framework/kv_cache/kv_cache.h"
+#include "framework/model/model_args.h"
+#include "framework/parallel_state/parallel_args.h"
+#include "framework/quant_args.h"
+#include "framework/state_dict/state_dict.h"
+#include "framework/state_dict/utils.h"
+#include "layers/common/linear.h"
+#include "layers/common/rms_norm_gated.h"
+
+namespace xllm {
+namespace layer {
+
+class Qwen3GatedDeltaNetBaseImpl : public torch::nn::Module {
+ public:
+  Qwen3GatedDeltaNetBaseImpl() = default;
+  Qwen3GatedDeltaNetBaseImpl(const ModelArgs& args,
+                             const QuantArgs& quant_args,
+                             const ParallelArgs& parallel_args,
+                             const torch::TensorOptions& options);
+
+  virtual void load_state_dict(const StateDict& state_dict) = 0;
+  virtual void verify_loaded_weights(const std::string& prefix) const = 0;
+
+  torch::Tensor forward(const torch::Tensor& hidden_states,
+                        const AttentionMetadata& attn_metadata,
+                        KVCache& kv_cache,
+                        const ModelInputParams& input_params);
+
+ protected:
+  virtual std::pair<torch::Tensor, torch::Tensor> project_decode_inputs(
+      const torch::Tensor& hidden_states) = 0;
+  virtual std::pair<torch::Tensor, torch::Tensor> project_flat_inputs(
+      const torch::Tensor& hidden_states) = 0;
+  // Qwen3.5 overrides this to project and reshape its separate qkv/z/b/a
+  // weights in every forward mode. Qwen3Next keeps qkvz/ba packed and returns
+  // nullopt to select the fused-split fallback.
+  virtual std::optional<
+      std::tuple<torch::Tensor, torch::Tensor, torch::Tensor, torch::Tensor>>
+  project_split_inputs(const torch::Tensor& hidden_states,
+                       const AttentionMetadata& attn_metadata) {
+    return std::nullopt;
+  }
+  virtual bool use_fla_ssm_state_layout() const { return false; }
+
+  void load_common_state_dict(const StateDict& state_dict);
+  void verify_common_loaded_weights(const std::string& prefix) const;
+
+  torch::Tensor get_linear_state_indices(const ModelInputParams& input_params,
+                                         const torch::Device& device) const;
+
+  std::pair<torch::Tensor, torch::Tensor> project_padded_inputs(
+      const torch::Tensor& hidden_states,
+      const AttentionMetadata& attn_metadata);
+
+  torch::Tensor reshape_qkvz_unpad(const AttentionMetadata& attn_metadata,
+                                   const torch::Tensor& padded_qkvz) const;
+
+  // Projection outputs are packed as [total_tokens, dim], while GDN kernels
+  // consume dense [batch, max_query_len, dim] tensors. Split the packed tokens
+  // by query length and pad each sequence before entering the kernels.
+  torch::Tensor reshape_projected_tokens_with_pad(
+      const AttentionMetadata& attn_metadata,
+      const torch::Tensor& projected_tokens) const;
+
+  std::tuple<torch::Tensor, torch::Tensor, torch::Tensor> process_mixed_qkv(
+      torch::Tensor& mixed_qkv) const;
+
+  int64_t num_k_heads_ = 0;
+  int64_t num_v_heads_ = 0;
+  int64_t head_k_dim_ = 0;
+  int64_t head_v_dim_ = 0;
+  int64_t k_size_ = 0;
+  int64_t v_size_ = 0;
+  int64_t tp_size_ = 1;
+  int64_t rank_ = 0;
+  int32_t conv_kernel_size_ = 0;
+
+  ColumnParallelLinear conv1d_{nullptr};
+  RowParallelLinear o_proj_{nullptr};
+  RmsNormGated norm_{nullptr};
+
+  DEFINE_WEIGHT(dt_bias);
+  DEFINE_WEIGHT(A_log);
+};
+
+}  // namespace layer
+}  // namespace xllm