sm90 still use row-major scale for permute-fp8

Author: zoooo0820

custom_ops/gpu_ops/moe/ep_moe_expert_dispatch.cu

@@ -942,9 +942,14 @@ __global__ void permute_x_fp8_kernel(
           }
 
         } else {
-          for (int s = tid; s < hidden_size_scale; s += blockDim.x) {
-            permute_scale[s * permute_scale_stride0 + dst_token_idx] =
-                scale[s * padded_num_rows + s_token_idx];
+          for (int v_id = tid; v_id < hidden_size_scale_int4;
+               v_id += blockDim.x) {
+            *(reinterpret_cast<int4*>(permute_scale +
+                                      dst_token_idx * hidden_size_scale) +
+              v_id) =
+                *(reinterpret_cast<const int4*>(scale + s_token_idx *
+                                                            hidden_size_scale) +
+                  v_id);
           }
         }
       }
@@ -1106,7 +1111,6 @@ std::vector<paddle::Tensor> EPMoeExpertDispatchFP8(
             m_indices};
   } else {
     permute_scale = GetEmptyTensor({token_nums_feed_to_ffn, hidden_size / 128},
-                                   {1, permute_scale_stride0},
                                    paddle::DataType::FLOAT32,
                                    place);
     EPMoeDispatchFP8Kernel<float>(input,

fastdeploy/model_executor/layers/moe/ep.py

@@ -173,7 +173,8 @@ def create_buffer(self):
                 self.num_nvl_bytes,
                 self.num_rdma_bytes,
                 low_latency_mode=True,
-                num_qps_per_rank=24,
+                # num_qps_per_rank=24,
+                num_qps_per_rank=48,
             )
             self.deepep_buffer.set_num_sms(14)  # TODO: tune in future
         else:
@@ -186,7 +187,8 @@ def create_buffer(self):
                     self.num_nvl_bytes,
                     self.num_rdma_bytes,
                     low_latency_mode=True,
-                    num_qps_per_rank=24,
+                    # num_qps_per_rank=24,
+                    num_qps_per_rank=48,
                 )
             else:
                 raise ValueError(f"Unknown generation phase: {self.moe_phase.phase}")

fastdeploy/model_executor/layers/moe/fused_moe_deepgemm_backend.py

@@ -99,6 +99,9 @@ def m_grouped_fp8_gemm_nt_contiguous_custom_python_op(
         (permute_input.shape[0], layer_added_weight_attrs_0.shape[1]),
         dtype=paddle.bfloat16,
     )
+    if disable_ue8m0_cast:
+        permute_scale = permute_scale.transpose([1, 0]).contiguous()
+        permute_scale = permute_scale.transpose([1, 0])
     # disable_ue8m0_cast is False for SM100
     m_grouped_fp8_gemm_nt_contiguous(
         (permute_input, permute_scale),
@@ -262,10 +265,14 @@ def apply_ep_prefill(
         x, x_scale_tensor = paddle.incubate.nn.functional.fp8_quant_blockwise(
             x,
             using_pow2_scale=self.quant_config.deepgemm_scale_ue8m0,
-            output_scale_transpose=True,
+            output_scale_transpose=self.quant_config.deepgemm_scale_ue8m0,
             using_ue8m0_scale=self.quant_config.deepgemm_scale_ue8m0,
         )
-        x_scale_tensor = x_scale_tensor.T[: x.shape[0]]
+        x_scale_tensor = (
+            x_scale_tensor[: x.shape[0]]
+            if not self.quant_config.deepgemm_scale_ue8m0
+            else x_scale_tensor.T[: x.shape[0]]
+        )
 
         event = deep_ep.Buffer.capture()
         let_another_thread_run()
@@ -502,10 +509,14 @@ def apply_tp(
         recv_x, recv_x_scale = paddle.incubate.nn.functional.fp8_quant_blockwise(
             x,
             using_pow2_scale=self.quant_config.deepgemm_scale_ue8m0,
-            output_scale_transpose=True,
+            output_scale_transpose=self.quant_config.deepgemm_scale_ue8m0,
             using_ue8m0_scale=self.quant_config.deepgemm_scale_ue8m0,
         )
-        recv_x_scale = recv_x_scale.T[: recv_x.shape[0]]
+        recv_x_scale = (
+            recv_x_scale[: recv_x.shape[0]]
+            if not self.quant_config.deepgemm_scale_ue8m0
+            else recv_x_scale.T[: recv_x.shape[0]]
+        )
         (
             permute_input,
             permute_scale,