change name

Author: GuoRen868

csrc/build_aclnn.sh

@@ -5,5 +5,5 @@ cd custom_ops/
 bash build.sh custom_ops -cascend910_93
 
 # install custom ops
-# ./output/CANN-custom_ops--linux.x86_64.run
-# export LD_LIBRARY_PATH=/usr/local/Ascend/ascend-toolkit/latest/opp/vendors/customize/op_api/lib/:${LD_LIBRARY_PATH}
+./build_out/custom_ops/run/CANN_ascend910_93_ubuntu_aarch64.run --install-path=/usr/local/Ascend/ascend-toolkit/latest/opp/
+source /usr/local/Ascend/ascend-toolkit/latest/opp/vendors/customize/bin/set_env.bash

csrc/custom_ops/kernels/fused_deep_moe/op_host/fused_deep_moe.cpp renamed to csrc/custom_ops/kernels/dispatch_gmm_combine_decode/op_host/dispatch_gmm_combine_decode.cpp

@@ -1,18 +1,18 @@
 /*
  * Copyright (c) Huawei Technologies Co., Ltd. 2025-2025. All rights reserved.
- * Description: FusedDeepMoe operator definition file
+ * Description: DispatchGmmCombineDecode operator definition file
  * Author: WANG Qiankun
  * Create: 2025-07-19
  * Note:
- * History: 2025-07-19 create FusedDeepMoe operator definition file
+ * History: 2025-07-19 create DispatchGmmCombineDecode operator definition file
  */
 #include "register/op_def_registry.h"
 
 namespace ops {
-class FusedDeepMoe : public OpDef
+class DispatchGmmCombineDecode : public OpDef
 {
 public:
-    explicit FusedDeepMoe(const char *name) : OpDef(name)
+    explicit DispatchGmmCombineDecode(const char *name) : OpDef(name)
     {
         this->Input("x")
             .ParamType(REQUIRED)
@@ -78,5 +78,5 @@ class FusedDeepMoe : public OpDef
     }
 };
 
-OP_ADD(FusedDeepMoe);
+OP_ADD(DispatchGmmCombineDecode);
 }  // namespace ops

csrc/custom_ops/kernels/fused_deep_moe/op_host/fused_deep_moe_infer.cpp renamed to csrc/custom_ops/kernels/dispatch_gmm_combine_decode/op_host/dispatch_gmm_combine_decode_infer.cpp

@@ -1,10 +1,10 @@
 /*
  * Copyright (c) Huawei Technologies Co., Ltd. 2025-2025. All rights reserved.
- * Description: FusedDeepMoe tiling function implementation file
+ * Description: DispatchGmmCombineDecode tiling function implementation file
  * Author: Guo Ren
  * Create: 2025-07-22
  * Note:
- * History: 2025-07-13 create FusedDeepMoe infer function file
+ * History: 2025-07-13 create DispatchGmmCombineDecode infer function file
  */
 
 #include <cstdint>
@@ -89,5 +89,5 @@ static ge::graphStatus InferDataType(gert::InferDataTypeContext *context)
     return ge::GRAPH_SUCCESS;
 }
 
-IMPL_OP(FusedDeepMoe).InferShape(InferShape).InferDataType(InferDataType);
+IMPL_OP(DispatchGmmCombineDecode).InferShape(InferShape).InferDataType(InferDataType);
 }  // namespace ge

csrc/custom_ops/kernels/fused_deep_moe/op_host/fused_deep_moe_tiling.cpp renamed to csrc/custom_ops/kernels/dispatch_gmm_combine_decode/op_host/dispatch_gmm_combine_decode_tiling.cpp

@@ -1,10 +1,10 @@
 /*
  * Copyright (c) Huawei Technologies Co., Ltd. 2025-2025. All rights reserved.
- * Description: FusedDeepMoe tiling function implementation file
+ * Description: DispatchGmmCombineDecode tiling function implementation file
  * Author: WANG Qiankun
  * Create: 2025-07-19
  * Note:
- * History: 2025-07-19 create FusedDeepMoe tiling function implementation file
+ * History: 2025-07-19 create DispatchGmmCombineDecode tiling function implementation file
  */
 #include <cstdio>
 #include <cstdint>
@@ -13,7 +13,7 @@
 #include "error_log.h"
 #include "graph/utils/type_utils.h"
 #include "register/op_def_registry.h"
-#include "../op_kernel/fused_deep_moe_tiling.h"
+#include "../op_kernel/dispatch_gmm_combine_decode_tiling.h"
 #include "tiling/platform/platform_ascendc.h"
 #include "tiling/hccl/hccl_tiling.h"
 
@@ -67,7 +67,7 @@ static size_t CeilUp(size_t x, size_t y)
 }
 
 static ge::graphStatus CheckTensorShape(gert::TilingContext *context, const char *nodeName,
-                                        FusedDeepMoeTilingData &tilingData)
+                                        DispatchGmmCombineDecodeTilingData &tilingData)
 {
     uint32_t epRankId = tilingData.disGmmDeqSwigluQuantGmmDeqComInfo.epRankId;
     uint32_t moeExpertNum = tilingData.disGmmDeqSwigluQuantGmmDeqComInfo.moeExpertNum;
@@ -127,7 +127,7 @@ static ge::graphStatus CheckTensorShape(gert::TilingContext *context, const char
     return ge::GRAPH_SUCCESS;
 }
 
-static ge::graphStatus CheckData(const char *nodeName, FusedDeepMoeTilingData &tilingData)
+static ge::graphStatus CheckData(const char *nodeName, DispatchGmmCombineDecodeTilingData &tilingData)
 {
     uint32_t batchSize = tilingData.disGmmDeqSwigluQuantGmmDeqComInfo.bs;
     OP_TILING_CHECK(batchSize < MIN_BATCH_SIZE, OP_LOGE(nodeName, "batchSize(bs) must >= %d.", MIN_BATCH_SIZE),
@@ -162,7 +162,7 @@ static ge::graphStatus CheckData(const char *nodeName, FusedDeepMoeTilingData &t
 }
 
 static ge::graphStatus GetAttrAndSetTilingData(gert::TilingContext *context, const char *nodeName,
-                                               FusedDeepMoeTilingData &tilingData, std::string &groupEp)
+                                               DispatchGmmCombineDecodeTilingData &tilingData, std::string &groupEp)
 {
     auto attrs = context->GetAttrs();
     OP_TILING_CHECK(attrs == nullptr, OP_LOGE(nodeName, "attrs is nullptr."), return ge::GRAPH_FAILED);
@@ -209,10 +209,10 @@ static ge::graphStatus GetAttrAndSetTilingData(gert::TilingContext *context, con
     return ge::GRAPH_SUCCESS;
 }
 
-static void SetHcommCfg(const gert::TilingContext *context, FusedDeepMoeTilingData *tiling, const std::string groupEp)
+static void SetHcommCfg(const gert::TilingContext *context, DispatchGmmCombineDecodeTilingData *tiling, const std::string groupEp)
 {
     const char *nodeName = context->GetNodeName();
-    OP_LOGD(nodeName, "FusedDeepMoe groupEp = %s", groupEp.c_str());
+    OP_LOGD(nodeName, "DispatchGmmCombineDecode groupEp = %s", groupEp.c_str());
     uint32_t opType = OP_TYPE_ALL_TO_ALL;
     std::string algConfigAllToAllStr = "AlltoAll=level0:fullmesh;level1:pairwise";
     std::string algConfigAllGatherStr = "AllGather=level0:ring";
@@ -223,7 +223,7 @@ static void SetHcommCfg(const gert::TilingContext *context, FusedDeepMoeTilingDa
 }
 
 static ge::graphStatus SetWorkSpace(gert::TilingContext *context, const char *nodeName,
-                                    FusedDeepMoeTilingData &tilingData)
+                                    DispatchGmmCombineDecodeTilingData &tilingData)
 {
     size_t *workSpaces = context->GetWorkspaceSizes(1);
     OP_TILING_CHECK(workSpaces == nullptr, OP_LOGE(nodeName, "workSpaces is nullptr."), return ge::GRAPH_FAILED);
@@ -263,10 +263,10 @@ static ge::graphStatus SetWorkSpace(gert::TilingContext *context, const char *no
     return ge::GRAPH_SUCCESS;
 }
 
-static ge::graphStatus FusedDeepMoeTilingFuncImpl(gert::TilingContext *context)
+static ge::graphStatus DispatchGmmCombineDecodeTilingFuncImpl(gert::TilingContext *context)
 {
     const char *nodeName = context->GetNodeName();
-    FusedDeepMoeTilingData *tilingData = context->GetTilingData<FusedDeepMoeTilingData>();
+    DispatchGmmCombineDecodeTilingData *tilingData = context->GetTilingData<DispatchGmmCombineDecodeTilingData>();
     OP_TILING_CHECK(tilingData == nullptr, OP_LOGE(nodeName, "tilingData is nullptr."), return ge::GRAPH_FAILED);
     std::string groupEp = "";
 
@@ -312,20 +312,20 @@ static ge::graphStatus FusedDeepMoeTilingFuncImpl(gert::TilingContext *context)
     return ge::GRAPH_SUCCESS;
 }
 
-static ge::graphStatus FusedDeepMoeTilingFunc(gert::TilingContext *context)
+static ge::graphStatus DispatchGmmCombineDecodeTilingFunc(gert::TilingContext *context)
 {
-    ge::graphStatus ret = FusedDeepMoeTilingFuncImpl(context);
+    ge::graphStatus ret = DispatchGmmCombineDecodeTilingFuncImpl(context);
     return ret;
 }
 
-struct FusedDeepMoeCompileInfo {};
-ge::graphStatus TilingParseForFusedDeepMoe(gert::TilingParseContext *context)
+struct DispatchGmmCombineDecodeCompileInfo {};
+ge::graphStatus TilingParseForDispatchGmmCombineDecode(gert::TilingParseContext *context)
 {
     (void)context;
     return ge::GRAPH_SUCCESS;
 }
 
-IMPL_OP_OPTILING(FusedDeepMoe)
-    .Tiling(FusedDeepMoeTilingFunc)
-    .TilingParse<FusedDeepMoeCompileInfo>(TilingParseForFusedDeepMoe);
+IMPL_OP_OPTILING(DispatchGmmCombineDecode)
+    .Tiling(DispatchGmmCombineDecodeTilingFunc)
+    .TilingParse<DispatchGmmCombineDecodeCompileInfo>(TilingParseForDispatchGmmCombineDecode);
 }  // namespace optiling

csrc/custom_ops/kernels/fused_deep_moe/op_kernel/fused_deep_moe.cpp renamed to csrc/custom_ops/kernels/dispatch_gmm_combine_decode/op_kernel/dispatch_gmm_combine_decode.cpp

@@ -1,16 +1,16 @@
 /*
  * Copyright (c) Huawei Technologies Co., Ltd. 2025-2025. All rights reserved.
- * Description: FusedDeepMoe operator kernel function implementation file
+ * Description: DispatchGmmCombineDecode operator kernel function implementation file
  * Author: WANG Qiankun
  * Create: 2025-07-19
  * Note:
- * History: 2025-07-19 create FusedDeepMoe operator kernel function implementation file
+ * History: 2025-07-19 create DispatchGmmCombineDecode operator kernel function implementation file
  */
-#include "fused_deep_moe.h"
+#include "dispatch_gmm_combine_decode.h"
 #include <kernel_operator.h>
 #include "lib/matmul_intf.h"
 
-extern "C" __global__ __aicore__ void fused_deep_moe(
+extern "C" __global__ __aicore__ void dispatch_gmm_combine_decode(
     // input
     GM_ADDR x, GM_ADDR expert_ids, GM_ADDR gmm1_permuted_weight, GM_ADDR gmm1_permuted_weight_scale,
     GM_ADDR gmm2_weight, GM_ADDR gmm2_weight_scale, GM_ADDR expert_smooth_scales, GM_ADDR expert_scales,
@@ -21,11 +21,11 @@ extern "C" __global__ __aicore__ void fused_deep_moe(
 {
     icache_preload(8);
     // New output recvCount
-    REGISTER_TILING_DEFAULT(FusedDeepMoeTilingData);
+    REGISTER_TILING_DEFAULT(DispatchGmmCombineDecodeTilingData);
     KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2);  // 1C2V
     GET_TILING_DATA(tiling_data, tiling);
     if constexpr (TILING_KEY_IS(0) || TILING_KEY_IS(1)) {
-        FusedDeepMoe<DTYPE_X, int32_t, false, TILING_KEY_VAR> op;
+        DispatchGmmCombineDecode<DTYPE_X, int32_t, false, TILING_KEY_VAR> op;
         op.Init(x, expert_ids, gmm1_permuted_weight, gmm1_permuted_weight_scale, gmm2_weight, gmm2_weight_scale,
                 expert_smooth_scales, expert_scales, output, outputRecvCount, workspace, nullptr, &tiling_data);
         op.Process();

csrc/custom_ops/kernels/fused_deep_moe/op_kernel/fused_deep_moe.h renamed to csrc/custom_ops/kernels/dispatch_gmm_combine_decode/op_kernel/dispatch_gmm_combine_decode.h

@@ -1,13 +1,13 @@
 /*
  * Copyright (c) Huawei Technologies Co., Ltd. 2025-2025. All rights reserved.
- * Description: FusedDeepMoe operator kernel function header file, for a3
+ * Description: DispatchGmmCombineDecode operator kernel function header file, for a3
  * Author: WANG Qiankun
  * Create: 2025-07-19
  * Note:
- * History: 2025-07-19 create FusedDeepMoe operator kernel function header file, for a3
+ * History: 2025-07-19 create DispatchGmmCombineDecode operator kernel function header file, for a3
  */
-#ifndef FUSED_DEEP_MOE_H
-#define FUSED_DEEP_MOE_H
+#ifndef DISPATCH_GMM_COMBINE_DECODE_H
+#define DISPATCH_GMM_COMBINE_DECODE_H
 
 #include "lib/matmul_intf.h"
 #include <kernel_operator.h>
@@ -29,8 +29,8 @@
 
 #include "operator/cam_moe_distribute_combine/op_kernel/a3/cam_moe_distribute_dispatch.h"
 
-#include "fused_deep_moe_tiling.h"
-#include "fused_deep_moe_base.h"
+#include "dispatch_gmm_combine_decode_tiling.h"
+#include "dispatch_gmm_combine_decode_base.h"
 
 #define ENABLE_GMM2_COMBINE
 
@@ -235,18 +235,18 @@ ACT_DEVICE void GmmDeq(GemmCoord problemShape, uint32_t groupCount, GM_ADDR gmGr
 }
 
 template <TemplateMC2TypeClass>
-class FusedDeepMoe
+class DispatchGmmCombineDecode
 {
 public:
-    __aicore__ inline FusedDeepMoe(){};
+    __aicore__ inline DispatchGmmCombineDecode(){};
     __aicore__ inline void Init(
         // input
         GM_ADDR x, GM_ADDR expert_ids, GM_ADDR gmm1_permuted_weight, GM_ADDR gmm1_permuted_weight_scale,
         GM_ADDR gmm2_weight, GM_ADDR gmm2_weight_scale, GM_ADDR expert_smooth_scales, GM_ADDR expert_scales,
         // output
         GM_ADDR output, GM_ADDR outputRecvCount,
         // system
-        GM_ADDR workspaceGM, AscendC::TPipe *pipe, const FusedDeepMoeTilingData *tilingData);
+        GM_ADDR workspaceGM, AscendC::TPipe *pipe, const DispatchGmmCombineDecodeTilingData *tilingData);
     __aicore__ inline void Process();
 
 private:
@@ -285,18 +285,18 @@ class FusedDeepMoe
 
     AscendC::TPipe *tpipe_{nullptr};
     __gm__ HcclOpResParam *winContext_{nullptr};
-    const FusedDeepMoeTilingData *tilingData_;
+    const DispatchGmmCombineDecodeTilingData *tilingData_;
 };
 
 template <TemplateMC2TypeClass>
-__aicore__ inline void FusedDeepMoe<TemplateMC2TypeFunc>::Init(
+__aicore__ inline void DispatchGmmCombineDecode<TemplateMC2TypeFunc>::Init(
     // input
     GM_ADDR x, GM_ADDR expert_ids, GM_ADDR gmm1_permuted_weight, GM_ADDR gmm1_permuted_weight_scale,
     GM_ADDR gmm2_weight, GM_ADDR gmm2_weight_scale, GM_ADDR expert_smooth_scales, GM_ADDR expert_scales,
     // output
     GM_ADDR output, GM_ADDR outputRecvCount,
     // system
-    GM_ADDR workspaceGM, AscendC::TPipe *pipe, const FusedDeepMoeTilingData *tilingData)
+    GM_ADDR workspaceGM, AscendC::TPipe *pipe, const DispatchGmmCombineDecodeTilingData *tilingData)
 {
     tpipe_ = pipe;
     blockDim_ = AscendC::GetBlockNum();
@@ -341,15 +341,15 @@ __aicore__ inline void FusedDeepMoe<TemplateMC2TypeFunc>::Init(
 }
 
 template <TemplateMC2TypeClass>
-__aicore__ inline void FusedDeepMoe<TemplateMC2TypeFunc>::Process()
+__aicore__ inline void DispatchGmmCombineDecode<TemplateMC2TypeFunc>::Process()
 {
 #ifdef ENABLE_GMM2_COMBINE
     if (g_coreType == AscendC::AIV) {
-        ((FusedDeepMoeTilingData *)tilingData_)->disGmmDeqSwigluQuantGmmDeqComInfo.aicNum = get_block_num();
+        ((DispatchGmmCombineDecodeTilingData *)tilingData_)->disGmmDeqSwigluQuantGmmDeqComInfo.aicNum = get_block_num();
         if constexpr (EXEC_FLAG & EXEC_FLAG_DEEP_FUSE) {
-            ((FusedDeepMoeTilingData *)tilingData_)->disGmmDeqSwigluQuantGmmDeqComInfo.aivNum = get_block_num();
+            ((DispatchGmmCombineDecodeTilingData *)tilingData_)->disGmmDeqSwigluQuantGmmDeqComInfo.aivNum = get_block_num();
         } else {
-            ((FusedDeepMoeTilingData *)tilingData_)->disGmmDeqSwigluQuantGmmDeqComInfo.aivNum =
+            ((DispatchGmmCombineDecodeTilingData *)tilingData_)->disGmmDeqSwigluQuantGmmDeqComInfo.aivNum =
                 get_block_num() * get_subblockdim();
         }
     }
@@ -444,4 +444,4 @@ __aicore__ inline void FusedDeepMoe<TemplateMC2TypeFunc>::Process()
                                layoutOutput, gmWorkspace, &combiner);
 #endif
 }
-#endif  // FUSED_DEEP_MOE_H
+#endif  // DISPATCH_GMM_COMBINE_DECODE_H

csrc/custom_ops/kernels/fused_deep_moe/op_kernel/fused_deep_moe_base.h renamed to csrc/custom_ops/kernels/dispatch_gmm_combine_decode/op_kernel/dispatch_gmm_combine_decode_base.h

@@ -6,12 +6,12 @@
  * Note:
  * History: 2025-07-19 Create a definition file for a distribution group related structure
  */
-#ifndef FUSED_DEEP_MOE_BASE_H
-#define FUSED_DEEP_MOE_BASE_H
+#ifndef DISPATCH_GMM_COMBINE_DECODE_BASE_H
+#define DISPATCH_GMM_COMBINE_DECODE_BASE_H
 
 #include "moe_distribute_base.h"
 
 #define TemplateMC2TypeClass typename ExpandXType, typename ExpandIdxType, bool IsNeedReduceScatter, uint32_t EXEC_FLAG
 #define TemplateMC2TypeFunc ExpandXType, ExpandIdxType, IsNeedReduceScatter, EXEC_FLAG
 
-#endif  // FUSED_DEEP_MOE_BASE_H
+#endif  // DISPATCH_GMM_COMBINE_DECODE_BASE_H

csrc/custom_ops/kernels/fused_deep_moe/op_kernel/fused_deep_moe_tiling.h renamed to csrc/custom_ops/kernels/dispatch_gmm_combine_decode/op_kernel/dispatch_gmm_combine_decode_tiling.h

@@ -1,18 +1,18 @@
 /*
  * Copyright (c) Huawei Technologies Co., Ltd. 2025-2025. All rights reserved.
- * Description: FusedDeepMoe tilingData definition file
+ * Description: DispatchGmmCombineDecode tilingData definition file
  * Author: WANG Qiankun
  * Create: 2025-07-19
  * Note:
- * History: 2025-07-19 create FusedDeepMoe tilingData definition file
+ * History: 2025-07-19 create DispatchGmmCombineDecode tilingData definition file
  */
 
-#ifndef FUSED_DEEP_MOE_TILING_H
-#define FUSED_DEEP_MOE_TILING_H
+#ifndef DISPATCH_GMM_COMBINE_DECODE_TILING_H
+#define DISPATCH_GMM_COMBINE_DECODE_TILING_H
 
 #include "kernel_tiling/kernel_tiling.h"
 
-struct FusedDeepMoeInfo {
+struct DispatchGmmCombineDecodeInfo {
     uint32_t epRankSize;           // epRankSize
     uint32_t epRankId;             // epRankId
     uint32_t moeExpertNum;         // moe expert number
@@ -31,10 +31,10 @@ struct FusedDeepMoeInfo {
     uint64_t gmm1HLen;
 };
 
-struct FusedDeepMoeTilingData {
+struct DispatchGmmCombineDecodeTilingData {
     Mc2InitTiling mc2InitTiling;
     Mc2CcTiling mc2CcTiling;
-    FusedDeepMoeInfo disGmmDeqSwigluQuantGmmDeqComInfo;
+    DispatchGmmCombineDecodeInfo disGmmDeqSwigluQuantGmmDeqComInfo;
 };
 
 constexpr uint32_t GM_ALIGN_BYTE = 512;
@@ -70,4 +70,4 @@ constexpr uint32_t WORKSPACE_STAGES = 4;
 
 constexpr uint32_t EXEC_FLAG_DEEP_FUSE = (1U << 0);
 
-#endif  // FUSED_DEEP_MOE_TILING_H
+#endif  // DISPATCH_GMM_COMBINE_DECODE_TILING_H