Star1108
2695
代码
代码
Issues206
Pull Requests376
Actions
项目讨论
Wiki
分析
Star1108
2695
已合并
Revert "modify quant_reduce_scatter hccl_context" #6541
yifuxiong创建于 6月8日
Revert "modify quant_reduce_scatter hccl_context" #6541
已合并
yifuxiong创建于 6月8日
从xiongyifu/ops-transformerpr_revert合入到cann/ops-transformermaster
共 37 个文件变更+283-1323
Mbuild.sh+0-4
@@ -1494,10 +1494,6 @@ if [ -n "${ascend_op_name}" ];then
1494 if [[ "${ascend_op_name}" == *"distribute_barrier"* ]] && [[ "${ascend_op_name}" != *"distribute_barrier_extend"* ]]; then1494 if [[ "${ascend_op_name}" == *"distribute_barrier"* ]] && [[ "${ascend_op_name}" != *"distribute_barrier_extend"* ]]; then
1495 ascend_op_name="${ascend_op_name};distribute_barrier_extend"1495 ascend_op_name="${ascend_op_name};distribute_barrier_extend"
1496 fi1496 fi
1497- # 编译quant_reduce_scatter的同时，把quant_reduce_scatter_v2也带上
1498- if [[ "${ascend_op_name}" == *"quant_reduce_scatter"* ]] && [[ "${ascend_op_name}" != *"quant_reduce_scatter_v2"* ]]; then
1499- ascend_op_name="${ascend_op_name};quant_reduce_scatter_v2"
1500- fi
1501 CUSTOM_OPTION="${CUSTOM_OPTION} -DASCEND_OP_NAME=${ascend_op_name}"1497 CUSTOM_OPTION="${CUSTOM_OPTION} -DASCEND_OP_NAME=${ascend_op_name}"
1502 if [[ "${ascend_op_name}" != *"fused_infer_attention_score"* ]] && [[ "${ascend_op_name}" != *"incre_flash_attention"* ]]; then1498 if [[ "${ascend_op_name}" != *"fused_infer_attention_score"* ]] && [[ "${ascend_op_name}" != *"incre_flash_attention"* ]]; then
1503 CUSTOM_OPTION="${CUSTOM_OPTION} -DENABLE_TILING_SINK=OFF"1499 CUSTOM_OPTION="${CUSTOM_OPTION} -DENABLE_TILING_SINK=OFF"
Mclassify_rule.yaml+0-5
@@ -639,11 +639,6 @@ mc2_infer@ops-transformer:
639 - ops/ops-transformer/mc2/quant_reduce_scatter/op_kernel/639 - ops/ops-transformer/mc2/quant_reduce_scatter/op_kernel/
640 - ops/ops-transformer/mc2/quant_reduce_scatter/CMakeLists.txt640 - ops/ops-transformer/mc2/quant_reduce_scatter/CMakeLists.txt
641 - ops/ops-transformer/mc2/quant_reduce_scatter/README.md641 - ops/ops-transformer/mc2/quant_reduce_scatter/README.md
642- - ops/ops-transformer/mc2/quant_reduce_scatter_v2/op_host/
643- - ops/ops-transformer/mc2/quant_reduce_scatter_v2/op_api/
644- - ops/ops-transformer/mc2/quant_reduce_scatter_v2/op_graph/
645- - ops/ops-transformer/mc2/quant_reduce_scatter_v2/op_kernel/
646- - ops/ops-transformer/mc2/quant_reduce_scatter_v2/CMakeLists.txt
647 - ops/ops-transformer/mc2/3rd/common/642 - ops/ops-transformer/mc2/3rd/common/
648 - ops/ops-transformer/mc2/3rd/ops_legacy/643 - ops/ops-transformer/mc2/3rd/ops_legacy/
649 - ops/ops-transformer/mc2/3rd/CMakeLists.txt644 - ops/ops-transformer/mc2/3rd/CMakeLists.txt
Mcmake/func.cmake+0-1
@@ -452,7 +452,6 @@ function(add_ops_src_copy)
452 "inplace_matmul_all_reduce_add_rms_norm;"452 "inplace_matmul_all_reduce_add_rms_norm;"
453 "quant_all_reduce;"453 "quant_all_reduce;"
454 "quant_reduce_scatter;"454 "quant_reduce_scatter;"
455- "quant_reduce_scatter_v2;"
456 "allto_all_matmul;"455 "allto_all_matmul;"
457 "matmul_allto_all;"456 "matmul_allto_all;"
458 "attention_to_ffn;"457 "attention_to_ffn;"
Mmc2/common/op_api/mc2_context.cpp+0-188
@@ -560,194 +560,6 @@ aclnnStatus Mc2Context::GetMc2RankSize(const char *groupEp, uint32_t &rankSize)
560 return ACLNN_SUCCESS;560 return ACLNN_SUCCESS;
561}561}
562 562 
563-/**
564- * @brief GetHcclCommResource for QuantReduceScatter
565- */
566-aclnnStatus Mc2Context::GetHcclCommResourceForQrs(const HcclComm &hcclHandle, const CommEngine &engine,
567- const CommProtocol &protocol,
568- Mc2QuantReduceScatterContext *mc2ContextStruct)
569-{
570- OP_LOGI("Start to get HCCL communication resource for QuantReduceScatter");
571- 
572- if (mc2ContextStruct->rankDim > HCCL_MTE_MAX_RANK_NUM) {
573- OP_LOGE(ACLNN_ERR_INNER,
574- "rankDim %u exceeds HCCL_MTE_MAX_RANK_NUM %u",
575- mc2ContextStruct->rankDim, HCCL_MTE_MAX_RANK_NUM);
576- return ACLNN_ERR_INNER;
577- }
578- 
579- uint32_t rankId = mc2ContextStruct->rankId;
580- std::vector<ChannelHandle> channels;
581- auto ret = GetHcclCommChannel(hcclHandle, mc2ContextStruct->rankDim, rankId, protocol, engine, channels);
582- if (ret != ACLNN_SUCCESS) {
583- return ret;
584- }
585- OP_LOGI("Get HCCL communication channel success, channel num is: %u", channels.size());
586- 
587- for (uint32_t i = 0; i < mc2ContextStruct->rankDim; ++i) {
588- void *tempBuffer = nullptr;
589- uint64_t bufSize = 0;
590- HcclResult hcclRet;
591- 
592- if (i == rankId) {
593- hcclRet = HcclGetHcclBuffer(hcclHandle, &tempBuffer, &hcclBuffSize_);
594- bufSize = hcclBuffSize_;
595- } else {
596- uint32_t idx = (i < rankId) ? i : (i - 1);
597- hcclRet = HcclChannelGetHcclBuffer(hcclHandle, channels[idx], &tempBuffer, &bufSize);
598- }
599- 
600- if (hcclRet != HCCL_SUCCESS || tempBuffer == nullptr) {
601- OP_LOGE(ACLNN_ERR_INNER, "Get HCCL buffer failed, src: %u, dst: %u", rankId, i);
602- return ACLNN_ERR_INNER;
603- }
604- 
605- mc2ContextStruct->windowsIn[i] = reinterpret_cast<uint64_t>(tempBuffer);
606- mc2ContextStruct->windowsOut[i] = reinterpret_cast<uint64_t>(static_cast<char *>(tempBuffer) + bufSize / 2);
607- }
608- 
609- OP_LOGI("Get HCCL CommResource for QuantReduceScatter success");
610- return ACLNN_SUCCESS;
611-}
612- 
613-/**
614- * @brief CreatMc2Context for QuantReduceScatter
615- */
616-aclnnStatus Mc2Context::CreatMc2ContextForQrs(const HcclComm &hcclHandle, const std::string &mc2ContextTag,
617- const CommEngine &engine, const CommProtocol &protocol,
618- Mc2QuantReduceScatterContext *mc2ContextStruct, void *&ctx,
619- uint64_t &hcclBuffSize)
620-{
621- OP_LOGI("Start to create HCCL context for QuantReduceScatter");
622- 
623- uint64_t ctxSize = sizeof(Mc2QuantReduceScatterContext);
624- auto hcclRet = HcclEngineCtxCreate(hcclHandle, mc2ContextTag.c_str(), engine, ctxSize, &ctx);
625- if (hcclRet != HCCL_SUCCESS) {
626- OP_LOGE(ACLNN_ERR_INNER, "Create HCCL context memory failed");
627- return ACLNN_ERR_INNER;
628- }
629- OP_LOGI("Create HCCL context for QuantReduceScatter success, context is: %p", ctx);
630- 
631- hcclRet = HcclGetRankId(hcclHandle, &mc2ContextStruct->rankId);
632- if (hcclRet != HCCL_SUCCESS) {
633- OP_LOGE(ACLNN_ERR_INNER, "Get rank ID failed");
634- return ACLNN_ERR_INNER;
635- }
636- OP_LOGI("Get rank ID success for QuantReduceScatter, rankId is: %u", mc2ContextStruct->rankId);
637- 
638- hcclRet = HcclGetRankSize(hcclHandle, &mc2ContextStruct->rankDim);
639- if (hcclRet != HCCL_SUCCESS) {
640- OP_LOGE(ACLNN_ERR_INNER, "Get rank size failed");
641- return ACLNN_ERR_INNER;
642- }
643- OP_LOGI("Get rank size for QuantReduceScatter success, rankSize is: %u", mc2ContextStruct->rankDim);
644- 
645- auto ret = GetHcclCommResourceForQrs(hcclHandle, engine, protocol, mc2ContextStruct);
646- if (ret != ACLNN_SUCCESS) {
647- OP_LOGE(ACLNN_ERR_INNER, "Get HCCL communication resource failed");
648- return ret;
649- }
650- 
651- mc2ContextStruct->workSpace = 0;
652- mc2ContextStruct->workSpaceSize = 0;
653- mc2ContextStruct->winSize = 0;
654- hcclRet = HcclEngineCtxCopy(hcclHandle, engine, mc2ContextTag.c_str(), mc2ContextStruct, ctxSize,
655- KOPY_DEFAULT_CTX_OFFSET);
656- if (hcclRet != HCCL_SUCCESS) {
657- OP_LOGE(ACLNN_ERR_INNER, "Copy context from host to device failed");
658- return ACLNN_ERR_INNER;
659- }
660- 
661- hcclBuffSize = hcclBuffSize_;
662- OP_LOGI("Copy context for QuantReduceScatter from host to device success");
663- return ACLNN_SUCCESS;
664-}
665- 
666-/**
667- * @brief CreatMc2ContextTensor for QuantReduceScatter
668- */
669-aclnnStatus Mc2Context::CreatMc2ContextTensorForQrs(void *ctx, aclTensor *&mc2Context)
670-{
671- OP_LOGI("Start to create Mc2Context Tensor for QuantReduceScatter");
672- 
673- if (ctx == nullptr) {
674- OP_LOGE(ACLNN_ERR_INNER, "Create Mc2Context Tensor failed, context is nullptr.");
675- return ACLNN_ERR_INNER;
676- }
677- 
678- uint64_t mc2ContextLength = sizeof(Mc2QuantReduceScatterContext);
679- int64_t shape[1] = {static_cast<int64_t>(mc2ContextLength / sizeof(uint32_t))};
680- int64_t strides[1] = {1};
681- 
682- mc2Context = aclCreateTensor(shape, 1, ACL_INT32, strides, 0, ACL_FORMAT_ND, shape, 1, ctx);
683- if (mc2Context == nullptr) {
684- OP_LOGE(ACLNN_ERR_INNER, "Create Mc2Context Tensor failed.");
685- return ACLNN_ERR_INNER;
686- }
687- 
688- OP_LOGI("CreatMc2ContextTensor for QuantReduceScatter Success");
689- return ACLNN_SUCCESS;
690-}
691- 
692-/**
693- * @brief GetMc2ContextTensor for QuantReduceScatter
694- */
695-aclnnStatus Mc2Context::GetMc2ContextTensorForQrs(const char *group, const char *opName, uint64_t &hcclBuffSize,
696- aclTensor *&mc2Context, int64_t &worldSize)
697-{
698- OP_LOGI("Start to get Mc2Context Tensor for QuantReduceScatter");
699- 
700- Mc2Context instance;
701- auto aclnnRet = instance.LoadHcclSymbols();
702- CHECK_RET(aclnnRet == ACLNN_SUCCESS, aclnnRet);
703- 
704- void *ctx = nullptr;
705- CommProtocol protocol;
706- std::string mc2ContextTag = std::string(group) + std::string(opName);
707- CommEngine engine = CommEngine::COMM_ENGINE_AIV;
708- hcclBuffSize = 0; // Default to 0, will be updated in CheckContextCache
709- 
710- aclnnRet = instance.ValidateContextTag(mc2ContextTag);
711- CHECK_RET(aclnnRet == ACLNN_SUCCESS, aclnnRet);
712- 
713- HcclComm hcclHandle;
714- aclnnRet = instance.GetCommHandle(group, hcclHandle);
715- CHECK_RET(aclnnRet == ACLNN_SUCCESS, aclnnRet);
716- 
717- uint32_t rankSize = 0;
718- auto hcclRet = instance.HcclGetRankSize(hcclHandle, &rankSize);
719- if (hcclRet != HCCL_SUCCESS) {
720- OP_LOGE(ACLNN_ERR_INNER, "Hccl get worldSize failed");
721- return ACLNN_ERR_INNER;
722- }
723- worldSize = rankSize;
724- OP_LOGI("Get worldSize success, worldSize is: %ld", worldSize);
725- 
726- aclnnRet = instance.CheckContextCache(hcclHandle, mc2ContextTag, engine, ctx, hcclBuffSize);
727- CHECK_RET(aclnnRet == ACLNN_SUCCESS, aclnnRet);
728- if (hcclBuffSize != 0) {
729- // Cache not found, need to create context
730- aclnnRet = instance.CreatMc2ContextTensorForQrs(ctx, mc2Context);
731- CHECK_RET(aclnnRet == ACLNN_SUCCESS, aclnnRet);
732- OP_LOGI("Found context cache, Get Mc2Context Tensor Success");
733- return ACLNN_SUCCESS;
734- }
735- 
736- aclnnRet = instance.GetCommProtocol(hcclHandle, protocol);
737- CHECK_RET(aclnnRet == ACLNN_SUCCESS, aclnnRet);
738- 
739- Mc2QuantReduceScatterContext mc2ContextStruct = {};
740- aclnnRet = instance.CreatMc2ContextForQrs(hcclHandle, mc2ContextTag, engine, protocol,
741- &mc2ContextStruct, ctx, hcclBuffSize);
742- CHECK_RET(aclnnRet == ACLNN_SUCCESS, aclnnRet);
743- 
744- aclnnRet = instance.CreatMc2ContextTensorForQrs(ctx, mc2Context);
745- CHECK_RET(aclnnRet == ACLNN_SUCCESS, aclnnRet);
746- 
747- OP_LOGI("Get Mc2QuantReduceScatterContext Tensor Success");
748- return ACLNN_SUCCESS;
749-}
750- 
751// 模板函数显式实例化563// 模板函数显式实例化
752template void *Mc2Context::GetHcclLibFunc<void *>(void *handle, const std::string &funcName);564template void *Mc2Context::GetHcclLibFunc<void *>(void *handle, const std::string &funcName);
753 565 
Mmc2/common/op_api/mc2_context.h+0-12
@@ -32,7 +32,6 @@
32#include "log/log.h"32#include "log/log.h"
33#include "opdev/common_types.h"33#include "opdev/common_types.h"
34#include "mc2_moe_context.h"34#include "mc2_moe_context.h"
35-#include "mc2_quant_reduce_scatter_context.h"
36 35 
37namespace Mc2Aclnn {36namespace Mc2Aclnn {
38 37 
@@ -41,8 +40,6 @@ public:
41 static aclnnStatus GetMc2ContextTensor(const char *groupEp, const char *opName, uint64_t &hcclBuffSize,40 static aclnnStatus GetMc2ContextTensor(const char *groupEp, const char *opName, uint64_t &hcclBuffSize,
42 aclTensor *&mc2Context);41 aclTensor *&mc2Context);
43 static aclnnStatus GetMc2RankSize(const char *groupEp, uint32_t &rankSize);42 static aclnnStatus GetMc2RankSize(const char *groupEp, uint32_t &rankSize);
44- static aclnnStatus GetMc2ContextTensorForQrs(const char *group, const char *opName, uint64_t &hcclBuffSize,
45- aclTensor *&mc2Context, int64_t &worldSize);
46 43 
47private:44private:
48 explicit Mc2Context();45 explicit Mc2Context();
@@ -72,15 +69,6 @@ private:
72 aclnnStatus CheckLinks(uint32_t &netLinkNum, CommLink *linksList);69 aclnnStatus CheckLinks(uint32_t &netLinkNum, CommLink *linksList);
73 aclnnStatus CheckContextCache(const HcclComm &hcclHandle, const std::string &mc2ContextTag,70 aclnnStatus CheckContextCache(const HcclComm &hcclHandle, const std::string &mc2ContextTag,
74 const CommEngine &engine, void *&ctx, uint64_t &hcclBuffSize);71 const CommEngine &engine, void *&ctx, uint64_t &hcclBuffSize);
75- /* for quant_reduce_scatter */
76- aclnnStatus GetHcclCommResourceForQrs(const HcclComm &hcclHandle, const CommEngine &engine,
77- const CommProtocol &protocol, Mc2QuantReduceScatterContext *mc2ContextStruct);
78- aclnnStatus CreatMc2ContextForQrs(const HcclComm &hcclHandle, const std::string &mc2ContextTag,
79- const CommEngine &engine, const CommProtocol &protocol,
80- Mc2QuantReduceScatterContext *mc2ContextStruct, void *&ctx,
81- uint64_t &hcclBuffSize);
82- aclnnStatus CreatMc2ContextTensorForQrs(void *ctx, aclTensor *&mc2Context);
83- 
84 const std::string GetLibPath();72 const std::string GetLibPath();
85 template <typename T>73 template <typename T>
86 T GetHcclLibFunc(void *handle, const std::string &funcName);74 T GetHcclLibFunc(void *handle, const std::string &funcName);
Dmc2/common/op_kernel/mc2_quant_reduce_scatter_context.h+0-36
@@ -1,36 +0,0 @@
1-/**
2- * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4- * CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9-**/
10- 
11-/*!
12- * \file mc2_quant_reduce_scatter_context.h
13- * \brief
14- */
15-#ifndef MC2_QUANT_REDUCE_SCATTER_CONTEXT_H
16-#define MC2_QUANT_REDUCE_SCATTER_CONTEXT_H
17- 
18-#include <cstdint>
19- 
20-namespace Mc2Aclnn {
21- 
22-constexpr uint32_t HCCL_MTE_MAX_RANK_NUM = 64;
23- 
24-struct Mc2QuantReduceScatterContext {
25- uint64_t workSpace; // client和server之间通信的地址
26- uint64_t workSpaceSize; // client和server之间通信的空间大小
27- uint32_t rankId; // 当前卡rankId
28- uint32_t rankDim; // 总卡数
29- uint64_t winSize; // ccu不使用
30- uint64_t windowsIn[HCCL_MTE_MAX_RANK_NUM]; // ccu不使用, MTE 数据区
31- uint64_t windowsOut[HCCL_MTE_MAX_RANK_NUM]; // ccu不使用，MTE 状态区
32-};
33- 
34-}
35- 
36-#endif // MC2_QUANT_REDUCE_SCATTER_CONTEXT_H
Mmc2/quant_all_reduce/op_host/CMakeLists.txt+1-1
@@ -18,7 +18,7 @@ if (BUILD_OPEN_PROJECT) # custom
18 OP_MC2_ENABLE ON 18 OP_MC2_ENABLE ON
19 OPTYPE quant_all_reduce quant_all_reduce ACLNNTYPE aclnn aclnn_inner)19 OPTYPE quant_all_reduce quant_all_reduce ACLNNTYPE aclnn aclnn_inner)
20 set(MC2_OPT ON PARENT_SCOPE)20 set(MC2_OPT ON PARENT_SCOPE)
21- set(quant_all_reduce_depends mc2/common mc2/3rd mc2/quant_reduce_scatter mc2/quant_reduce_scatter_v2 PARENT_SCOPE)21+ set(quant_all_reduce_depends mc2/common mc2/3rd mc2/quant_reduce_scatter PARENT_SCOPE)
22 set(SUB_MC2_COMPILE TRUE PARENT_SCOPE)22 set(SUB_MC2_COMPILE TRUE PARENT_SCOPE)
23 23 
24 # --cce-auto-sync=off：指定CCE编译器是否自动执行线程间或模块间的同步操作 24 # --cce-auto-sync=off：指定CCE编译器是否自动执行线程间或模块间的同步操作
Mmc2/quant_all_reduce/op_host/op_tiling/quant_all_reduce_tiling.cpp+12-20
@@ -67,8 +67,7 @@ static ge::graphStatus SetHcommCfg(const gert::TilingContext *context, QuantAllR
67 * @param tilingData: 框架根据context的opName匹配tiling模板，计算产生的tilingData67 * @param tilingData: 框架根据context的opName匹配tiling模板，计算产生的tilingData
68 * @return68 * @return
69 */69 */
70-static void SetTilingData(gert::TilingContext *context, QuantAllReduceTilingData &tilingData,70+static void SetTilingData(gert::TilingContext *context, QuantAllReduceTilingData &tilingData)
71- const QuantReduceScatterConfig& config)
72{71{
73 fe::PlatFormInfos *platformInfoPtr = context->GetPlatformInfo();72 fe::PlatFormInfos *platformInfoPtr = context->GetPlatformInfo();
74 platform_ascendc::PlatformAscendC ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);73 platform_ascendc::PlatformAscendC ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
@@ -76,25 +75,22 @@ static void SetTilingData(gert::TilingContext *context, QuantAllReduceTilingData
76 uint32_t aivNum = ascendcPlatform.GetCoreNumAiv();75 uint32_t aivNum = ascendcPlatform.GetCoreNumAiv();
77 context->SetBlockDim(ascendcPlatform.CalcTschBlockDim(aivNum, 0, aivNum));76 context->SetBlockDim(ascendcPlatform.CalcTschBlockDim(aivNum, 0, aivNum));
78 tilingData.quantAllReduceTilingInfo.aivNum = aivNum;77 tilingData.quantAllReduceTilingInfo.aivNum = aivNum;
79- uint64_t xValueBS = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);78+ uint64_t xValueBS = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);
80- uint64_t xValueH = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);79+ uint64_t xValueH = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
81- uint64_t scalesValueH = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_ONE);80+ uint64_t scalesValueH = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(DIM_ONE);
82 // context->GetInputShape在函数CheckInputTensorDim中已经校验81 // context->GetInputShape在函数CheckInputTensorDim中已经校验
83- if (context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDimNum() == THREE_DIMS) {82+ if (context->GetInputShape(X_INDEX)->GetStorageShape().GetDimNum() == THREE_DIMS) {
84- xValueBS = xValueBS * context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);83+ xValueBS = xValueBS * context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
85- xValueH = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_TWO);84+ xValueH = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_TWO);
86- scalesValueH = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_TWO);85+ scalesValueH = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(DIM_TWO);
87 }86 }
88 tilingData.quantAllReduceTilingInfo.bs = xValueBS;87 tilingData.quantAllReduceTilingInfo.bs = xValueBS;
89 tilingData.quantAllReduceTilingInfo.hiddenSize = xValueH;88 tilingData.quantAllReduceTilingInfo.hiddenSize = xValueH;
90 tilingData.quantAllReduceTilingInfo.scaleHiddenSize = scalesValueH;89 tilingData.quantAllReduceTilingInfo.scaleHiddenSize = scalesValueH;
91 tilingData.quantAllReduceTilingInfo.totalWinSize = mc2tiling::Mc2TilingUtils::GetMaxWindowSize();90 tilingData.quantAllReduceTilingInfo.totalWinSize = mc2tiling::Mc2TilingUtils::GetMaxWindowSize();
92- tilingData.quantAllReduceTilingInfo.isMc2Context = config.isMc2Context;
93}91}
94 92 
95-/**93+// 基于 TARGET_ITER 公式计算 host 推荐的 xPerBlock，写入 tilingData
96- * @brief 基于 TARGET_ITER 公式计算 host 推荐的 xPerBlock，写入 tilingData
97- */
98static void SetXPerBlock(QuantAllReduceTilingData &tilingData)94static void SetXPerBlock(QuantAllReduceTilingData &tilingData)
99{95{
100 constexpr uint32_t TARGET_ITER = 3U; // T=3 命中 DoubleBuffer 甜点96 constexpr uint32_t TARGET_ITER = 3U; // T=3 命中 DoubleBuffer 甜点
@@ -145,16 +141,12 @@ static ge::graphStatus QuantAllReduceTilingFunc(gert::TilingContext *context)
145 OP_TILING_CHECK(QuantReduceScatterUtilTiling::CheckNpuArch(context) != ge::GRAPH_SUCCESS,141 OP_TILING_CHECK(QuantReduceScatterUtilTiling::CheckNpuArch(context) != ge::GRAPH_SUCCESS,
146 OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(nodeName, "npuArch", "non-DAV_3510", "The value of npuArch must be DAV_3510"),142 OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(nodeName, "npuArch", "non-DAV_3510", "The value of npuArch must be DAV_3510"),
147 return ge::GRAPH_FAILED);143 return ge::GRAPH_FAILED);
148- QuantReduceScatterConfig config;144+ OP_TILING_CHECK(QuantReduceScatterUtilTiling::CheckTilingFunc(context, runInfo, OpType::OP_QUANT_ALL_REDUCE) !=
149- config.X_INDEX = 0;145+ ge::GRAPH_SUCCESS,
150- config.SCALES_INDEX = 1;
151- config.isMc2Context = false;
152- OP_TILING_CHECK(QuantReduceScatterUtilTiling::CheckTilingFunc(context, runInfo,
153- OpType::OP_QUANT_ALL_REDUCE, config) != ge::GRAPH_SUCCESS,
154 OP_LOGE(nodeName, "tiling check failed in quant_all_reduce."), return ge::GRAPH_FAILED);146 OP_LOGE(nodeName, "tiling check failed in quant_all_reduce."), return ge::GRAPH_FAILED);
155 OP_TILING_CHECK(SetHcommCfg(context, tilingData, runInfo) != ge::GRAPH_SUCCESS,147 OP_TILING_CHECK(SetHcommCfg(context, tilingData, runInfo) != ge::GRAPH_SUCCESS,
156 OP_LOGE(nodeName, "SetHCommCfg failed."), return ge::GRAPH_FAILED);148 OP_LOGE(nodeName, "SetHCommCfg failed."), return ge::GRAPH_FAILED);
157- SetTilingData(context, *tilingData, config);149+ SetTilingData(context, *tilingData);
158 SetXPerBlock(*tilingData);150 SetXPerBlock(*tilingData);
159 SetTilingKey(context);151 SetTilingKey(context);
160 PrintTilingDataInfo(context, *tilingData);152 PrintTilingDataInfo(context, *tilingData);
Mmc2/quant_all_reduce/op_kernel/quant_all_reduce_tiling_data.h+0-1
@@ -27,7 +27,6 @@ struct QuantAllReduceTilingInfo {
27 uint64_t totalWinSize; // Win区总大小，即HCCL_BUFFER_SIZE27 uint64_t totalWinSize; // Win区总大小，即HCCL_BUFFER_SIZE
28 uint32_t xPerBlock; // host 侧基于 TARGET_ITER 公式推荐的每块元素数28 uint32_t xPerBlock; // host 侧基于 TARGET_ITER 公式推荐的每块元素数
29 uint32_t alignBlock; // xPerBlock 对齐粒度（元素数，host/kernel共享）29 uint32_t alignBlock; // xPerBlock 对齐粒度（元素数，host/kernel共享）
30- bool isMc2Context;
31};30};
32 31 
33struct QuantAllReduceTilingData {32struct QuantAllReduceTilingData {
Mmc2/quant_reduce_scatter/op_api/aclnn_quant_reduce_scatter.cpp+135-12
@@ -13,7 +13,6 @@
13 * \brief13 * \brief
14 */14 */
15#include "aclnn_quant_reduce_scatter.h"15#include "aclnn_quant_reduce_scatter.h"
16-#include "aclnn_quant_reduce_scatter_base.h"
17#include "securec.h"16#include "securec.h"
18#include "acl/acl.h"17#include "acl/acl.h"
19#include "common/utils/op_mc2.h"18#include "common/utils/op_mc2.h"
@@ -30,19 +29,143 @@
30 29 
31using namespace op;30using namespace op;
32 31 
33-extern "C" aclnnStatus aclnnQuantReduceScatterGetWorkspaceSize(const aclTensor* x, const aclTensor* scales,32+namespace {
34- const char* group, const char* reduceOp,33+enum class NnopbaseHcclServerType : uint32_t {
35- aclTensor* output,34+ NNOPBASE_HCCL_SERVER_TYPE_AICPU = 0,
36- uint64_t* workspaceSize, aclOpExecutor** executor)35+ NNOPBASE_HCCL_SERVER_TYPE_MTE,
36+ NNOPBASE_HCCL_SERVER_TYPE_CCU,
37+ NNOPBASE_HCCL_SERVER_TYPE_END
38+};
39+ 
40+static constexpr size_t HCCL_GROUP_NAME_LENGTH_MAX = 128U; // group长度小于128字符
41+ 
42+// 根据API定义，列出K-G量化所能支持的所有dtype
43+const std::initializer_list<op::DataType> X_DTYPE_KG_SUPPORT_LIST = {
44+ op::DataType::DT_INT8, op::DataType::DT_HIFLOAT8, op::DataType::DT_FLOAT8_E4M3FN,
45+ op::DataType::DT_FLOAT8_E5M2
46+};
47+const std::initializer_list<op::DataType> SCALES_DTYPE_KG_SUPPORT_LIST = {
48+ op::DataType::DT_FLOAT
49+};
50+ 
51+// 根据API定义，列出MX量化所能支持的所有dtype
52+const std::initializer_list<op::DataType> X_DTYPE_MX_SUPPORT_LIST = {
53+ op::DataType::DT_FLOAT8_E4M3FN, op::DataType::DT_FLOAT8_E5M2
54+};
55+const std::initializer_list<op::DataType> SCALES_DTYPE_MX_SUPPORT_LIST = {
56+ op::DataType::DT_FLOAT8_E8M0
57+};
58+ 
59+const std::initializer_list<op::DataType> OUTPUT_DTYPE_SUPPORT_LIST = {
60+ op::DataType::DT_FLOAT16, op::DataType::DT_BF16, op::DataType::DT_FLOAT
61+};
62+ 
63+// 检查入参是否为nullptr
64+static bool CheckNotNull(const aclTensor* x, const aclTensor* scales, const aclTensor* output)
37{65{
38- OP_LOGD("aclnnQuantReduceScatterGetWorkspaceSize start");66+ OP_CHECK_NULL(x, return false);
39- return aclnnQuantReduceScatterBaseGetWorkspaceSize(x, scales, group, reduceOp, output,67+ OP_CHECK_NULL(scales, return false);
40- workspaceSize, executor);68+ OP_CHECK_NULL(output, return false);
69+ return true;
41}70}
42 71 
43-extern "C" aclnnStatus aclnnQuantReduceScatter(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,72+// 检查x、scales、output的数据类型是否在算子的支持列表之内
44- const aclrtStream stream)73+static bool CheckKGAllDtypesValid(const aclTensor* x, const aclTensor* scales, const aclTensor* output)
45{74{
46- OP_LOGD("aclnnQuantReduceScatter start");75+ if (CheckType(x->GetDataType(), X_DTYPE_KG_SUPPORT_LIST) && CheckType(scales->GetDataType(), SCALES_DTYPE_KG_SUPPORT_LIST) &&
47- return aclnnQuantReduceScatterBase(workspace, workspaceSize, executor, stream);76+ CheckType(output->GetDataType(), OUTPUT_DTYPE_SUPPORT_LIST)) {
77+ return true;
78+ } else {
79+ return false;
80+ }
81+}
82+ 
83+static bool CheckMXAllDtypesValid(const aclTensor* x, const aclTensor* scales, const aclTensor* output)
84+{
85+ if (CheckType(x->GetDataType(), X_DTYPE_MX_SUPPORT_LIST) && CheckType(scales->GetDataType(), SCALES_DTYPE_MX_SUPPORT_LIST) &&
86+ CheckType(output->GetDataType(), OUTPUT_DTYPE_SUPPORT_LIST)) {
87+ return true;
88+ } else {
89+ return false;
90+ }
91+}
92+ 
93+static bool CheckAllDtypesValid(const aclTensor* x, const aclTensor* scales, const aclTensor* output)
94+{
95+ bool isAllDtypesValid = false;
96+ isAllDtypesValid = CheckKGAllDtypesValid(x, scales, output) || CheckMXAllDtypesValid(x, scales, output);
97+ if (!isAllDtypesValid) {
98+ OP_LOGE_FOR_INVALID_DTYPE_WITH_REASON("aclnnQuantReduceScatter", "x/scales/output",
99+ (std::string(op::ToString(x->GetDataType()).GetString()) + "/" +
100+ op::ToString(scales->GetDataType()).GetString() + "/" +
101+ op::ToString(output->GetDataType()).GetString()).c_str(),
102+ "The dtypes of x, scales and output must be valid");
103+ }
104+ return isAllDtypesValid;
105+}
106+ 
107+static bool CheckGroupLength(const char* group)
108+{
109+ if (group == nullptr) {
110+ OP_LOGE_WITH_INVALID_INPUT("aclnnQuantReduceScatter", "group");
111+ return false;
112+ }
113+ 
114+ size_t groupLen = strnlen(group, HCCL_GROUP_NAME_LENGTH_MAX); // group长度≥128字符, 返回HCCL_GROUP_NAME_LENGTH_MAX
115+ if (groupLen >= HCCL_GROUP_NAME_LENGTH_MAX) {
116+ OP_LOGE_FOR_INVALID_VALUE_WITH_REASON("aclnnQuantReduceScatter", "group",
117+ "length exceeds " + std::to_string(HCCL_GROUP_NAME_LENGTH_MAX),
118+ "The length of group must be less than " + std::to_string(HCCL_GROUP_NAME_LENGTH_MAX) + " characters");
119+ return false;
120+ }
121+ 
122+ return true;
123+}
124+ 
125+static aclnnStatus CheckParams(const aclTensor* x, const aclTensor* scales, const char* group, const aclTensor* output)
126+{
127+ // 1. 检查参数是否为空指针
128+ CHECK_RET(CheckNotNull(x, scales, output), ACLNN_ERR_PARAM_NULLPTR);
129+ // 2. 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验
130+ CHECK_RET(CheckAllDtypesValid(x, scales, output), ACLNN_ERR_PARAM_INVALID);
131+ // 3. 检查group参数是否在要求范围之内
132+ CHECK_RET(CheckGroupLength(group), ACLNN_ERR_PARAM_INVALID);
133+ 
134+ return ACLNN_SUCCESS;
135+}
136+}
137+ 
138+extern "C" aclnnStatus aclnnInnerQuantReduceScatterGetWorkspaceSize(const aclTensor* x, const aclTensor* scales,
139+ const char* group, const char* reduceOp,
140+ uint64_t yDtype, int64_t worldSize, aclTensor* output,
141+ uint64_t* workspaceSize, aclOpExecutor** executor);
142+extern "C" aclnnStatus aclnnInnerQuantReduceScatter(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
143+ const aclrtStream stream);
144+extern "C" void __attribute__((weak)) NnopbaseSetHcclServerType(void *executor, NnopbaseHcclServerType sType);
145+ 
146+extern "C" aclnnStatus aclnnQuantReduceScatterGetWorkspaceSize(const aclTensor* x, const aclTensor* scales, const char* group,
147+ const char* reduceOp, aclTensor* output, uint64_t* workspaceSize,
148+ aclOpExecutor** executor)
149+{
150+ aclnnStatus retParam = CheckParams(x, scales, group, output);
151+ CHECK_RET(retParam == ACLNN_SUCCESS, retParam);
152+ uint64_t yDtype = static_cast<uint64_t>(output->GetDataType());
153+ int64_t worldSize = -1;
154+ aclnnStatus ret = aclnnInnerQuantReduceScatterGetWorkspaceSize(x, scales, const_cast<char*>(group),
155+ const_cast<char*>(reduceOp), yDtype, worldSize, output, workspaceSize, executor);
156+ OP_LOGD("QuantReduceScatter, aclnnnGetWorkspaceSize ret %d.", ret);
157+ return ret;
158+}
159+ 
160+extern "C" aclnnStatus aclnnQuantReduceScatter(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, const aclrtStream stream)
161+{
162+ if (NnopbaseSetHcclServerType) {
163+ NnopbaseSetHcclServerType(executor, NnopbaseHcclServerType::NNOPBASE_HCCL_SERVER_TYPE_MTE);
164+ }
165+ aclnnStatus ret = aclnnInnerQuantReduceScatter(workspace, workspaceSize, executor, stream);
166+ if (ret != ACLNN_SUCCESS) {
167+ OP_LOGE_LIBOPAPI_REPORT("aclnnQuantReduceScatter", "This is an error in launch aicore");
168+ return ACLNN_ERR_INNER;
169+ }
170+ return ACLNN_SUCCESS;
48}171}
Dmc2/quant_reduce_scatter/op_api/aclnn_quant_reduce_scatter_base.cpp+0-225
@@ -1,225 +0,0 @@
1-/**
2- * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4- * CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9-**/
10- 
11-/*!
12- * \file aclnn_quant_reduce_scatter_base.cpp
13- * \brief
14- */
15-#include "aclnn_quant_reduce_scatter_base.h"
16-#include "securec.h"
17-#include "acl/acl.h"
18-#include "common/utils/op_mc2.h"
19-#include "common/utils/op_mc2_def.h"
20-#include "aclnn_kernels/common/op_error_check.h"
21-#include "opdev/common_types.h"
22-#include "opdev/make_op_executor.h"
23-#include "opdev/op_dfx.h"
24-#include "opdev/op_executor.h"
25-#include "opdev/op_log.h"
26-#include "opdev/platform.h"
27-#include "common/utils/hccl_util.h"
28-#include "log/log.h"
29-#include "mc2/common/op_kernel/mc2_quant_reduce_scatter_context.h"
30- 
31-#ifdef BUILD_OPEN_PROJECT
32- 
33-#include "version/hcomm_version.h"
34-#define HCCL_CHANNEL_SUPPORT_VERSION 89999700
35- 
36-#if HCOMM_VERSION_NUM >= HCCL_CHANNEL_SUPPORT_VERSION
37-#include "common/op_api/mc2_context.h"
38-#endif // HCOMM_VERSION_NUM >= HCCL_CHANNEL_SUPPORT_VERSION
39- 
40-#endif // BUILD_OPEN_PROJECT
41- 
42-namespace {
43- 
44-using namespace op;
45- 
46-enum class NnopbaseHcclServerType : uint32_t {
47- NNOPBASE_HCCL_SERVER_TYPE_AICPU = 0,
48- NNOPBASE_HCCL_SERVER_TYPE_MTE,
49- NNOPBASE_HCCL_SERVER_TYPE_CCU,
50- NNOPBASE_HCCL_SERVER_TYPE_END
51-};
52- 
53-static constexpr size_t HCCL_GROUP_NAME_LENGTH_MAX = 128U; // group长度小于128字符
54- 
55-// 根据API定义，列出K-G量化所能支持的所有dtype
56-const std::initializer_list<op::DataType> X_DTYPE_KG_SUPPORT_LIST = {
57- op::DataType::DT_INT8, op::DataType::DT_HIFLOAT8, op::DataType::DT_FLOAT8_E4M3FN,
58- op::DataType::DT_FLOAT8_E5M2
59-};
60-const std::initializer_list<op::DataType> SCALES_DTYPE_KG_SUPPORT_LIST = {
61- op::DataType::DT_FLOAT
62-};
63- 
64-// 根据API定义，列出MX量化所能支持的所有dtype
65-const std::initializer_list<op::DataType> X_DTYPE_MX_SUPPORT_LIST = {
66- op::DataType::DT_FLOAT8_E4M3FN, op::DataType::DT_FLOAT8_E5M2
67-};
68-const std::initializer_list<op::DataType> SCALES_DTYPE_MX_SUPPORT_LIST = {
69- op::DataType::DT_FLOAT8_E8M0
70-};
71- 
72-const std::initializer_list<op::DataType> OUTPUT_DTYPE_SUPPORT_LIST = {
73- op::DataType::DT_FLOAT16, op::DataType::DT_BF16, op::DataType::DT_FLOAT
74-};
75- 
76-// 检查入参是否为nullptr
77-static bool CheckNotNull(const aclTensor* x, const aclTensor* scales, const aclTensor* output)
78-{
79- OP_CHECK_NULL(x, return false);
80- OP_CHECK_NULL(scales, return false);
81- OP_CHECK_NULL(output, return false);
82- return true;
83-}
84- 
85-// 检查x、scales、output的数据类型是否在算子的支持列表之内
86-static bool CheckKGAllDtypesValid(const aclTensor* x, const aclTensor* scales, const aclTensor* output)
87-{
88- if (CheckType(x->GetDataType(), X_DTYPE_KG_SUPPORT_LIST) &&
89- CheckType(scales->GetDataType(), SCALES_DTYPE_KG_SUPPORT_LIST) &&
90- CheckType(output->GetDataType(), OUTPUT_DTYPE_SUPPORT_LIST)) {
91- return true;
92- } else {
93- return false;
94- }
95-}
96- 
97-static bool CheckMXAllDtypesValid(const aclTensor* x, const aclTensor* scales, const aclTensor* output)
98-{
99- if (CheckType(x->GetDataType(), X_DTYPE_MX_SUPPORT_LIST) &&
100- CheckType(scales->GetDataType(), SCALES_DTYPE_MX_SUPPORT_LIST) &&
101- CheckType(output->GetDataType(), OUTPUT_DTYPE_SUPPORT_LIST)) {
102- return true;
103- } else {
104- return false;
105- }
106-}
107- 
108-static bool CheckAllDtypesValid(const aclTensor* x, const aclTensor* scales, const aclTensor* output)
109-{
110- bool isAllDtypesValid = false;
111- isAllDtypesValid = CheckKGAllDtypesValid(x, scales, output) || CheckMXAllDtypesValid(x, scales, output);
112- if (!isAllDtypesValid) {
113- OP_LOGE_FOR_INVALID_DTYPE_WITH_REASON("aclnnQuantReduceScatter", "x/scales/output",
114- (std::string(op::ToString(x->GetDataType()).GetString()) + "/" +
115- op::ToString(scales->GetDataType()).GetString() + "/" +
116- op::ToString(output->GetDataType()).GetString()).c_str(),
117- "Tensors x, scales and output are not simultaneously supported");
118- }
119- return isAllDtypesValid;
120-}
121- 
122-static bool CheckGroupLength(const char* group)
123-{
124- if (group == nullptr) {
125- OP_LOGE_WITH_INVALID_INPUT("aclnnQuantReduceScatter", "group");
126- return false;
127- }
128- 
129- size_t groupLen = strnlen(group, HCCL_GROUP_NAME_LENGTH_MAX); // group长度≥128字符, 返回HCCL_GROUP_NAME_LENGTH_MAX
130- if (groupLen >= HCCL_GROUP_NAME_LENGTH_MAX) {
131- OP_LOGE_FOR_INVALID_VALUE_WITH_REASON("aclnnQuantReduceScatter", "group",
132- "length exceeds " + std::to_string(HCCL_GROUP_NAME_LENGTH_MAX),
133- "Limit the length of the group to less than " + std::to_string(HCCL_GROUP_NAME_LENGTH_MAX) + " characters");
134- return false;
135- }
136- 
137- return true;
138-}
139- 
140-static aclnnStatus CheckParams(const aclTensor* x, const aclTensor* scales, const char* group, const aclTensor* output)
141-{
142- // 1. 检查参数是否为空指针
143- CHECK_RET(CheckNotNull(x, scales, output), ACLNN_ERR_PARAM_NULLPTR);
144- // 2. 检查输入的数据类型是否在API支持的数据类型范围之内，需要根据api定义校验
145- CHECK_RET(CheckAllDtypesValid(x, scales, output), ACLNN_ERR_PARAM_INVALID);
146- // 3. 检查group参数是否在要求范围之内
147- CHECK_RET(CheckGroupLength(group), ACLNN_ERR_PARAM_INVALID);
148- 
149- return ACLNN_SUCCESS;
150-}
151- 
152-} // namespace
153- 
154-extern "C" void __attribute__((weak)) NnopbaseSetHcclServerType(void *executor, NnopbaseHcclServerType sType);
155- 
156-// 走aclnn_quant_reduce_scatter_v2
157-#if defined(BUILD_OPEN_PROJECT) && HCOMM_VERSION_NUM >= HCCL_CHANNEL_SUPPORT_VERSION
158- 
159-extern "C" aclnnStatus aclnnInnerQuantReduceScatterV2GetWorkspaceSize(const aclTensor *context, const aclTensor* x,
160- const aclTensor* scales,
161- int64_t hcclBufferSize,
162- const char* reduceOp, int64_t yDtype,
163- int64_t worldSize, aclTensor* output,
164- uint64_t* workspaceSize,
165- aclOpExecutor** executor);
166-extern "C" aclnnStatus aclnnInnerQuantReduceScatterV2(void* workspace, uint64_t workspaceSize,
167- aclOpExecutor* executor, const aclrtStream stream);
168- 
169-#else
170- 
171-extern "C" aclnnStatus aclnnInnerQuantReduceScatterGetWorkspaceSize(const aclTensor* x, const aclTensor* scales,
172- const char* group, const char* reduceOp,
173- int64_t yDtype, int64_t worldSize,
174- aclTensor* output, uint64_t* workspaceSize,
175- aclOpExecutor** executor);
176-extern "C" aclnnStatus aclnnInnerQuantReduceScatter(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
177- const aclrtStream stream);
178-#endif
179- 
180-extern "C" aclnnStatus aclnnQuantReduceScatterBaseGetWorkspaceSize(const aclTensor* x, const aclTensor* scales,
181- const char* group, const char* reduceOp,
182- aclTensor* output, uint64_t* workspaceSize,
183- aclOpExecutor** executor)
184-{
185- aclnnStatus retParam = CheckParams(x, scales, group, output);
186- CHECK_RET(retParam == ACLNN_SUCCESS, retParam);
187- uint64_t yDtype = static_cast<uint64_t>(output->GetDataType());
188- int64_t worldSize = -1;
189- 
190- aclnnStatus ret = ACL_SUCCESS;
191-#if defined(BUILD_OPEN_PROJECT) && HCOMM_VERSION_NUM >= HCCL_CHANNEL_SUPPORT_VERSION
192- aclTensor *mc2Context = nullptr;
193- uint64_t hcclBuffSize = 0;
194- const char *opName = "quant_reduce_scatter_v2";
195- auto aclnnRet = Mc2Aclnn::Mc2Context::GetMc2ContextTensorForQrs(group, opName, hcclBuffSize,
196- mc2Context, worldSize);
197- CHECK_RET(aclnnRet == ACLNN_SUCCESS, aclnnRet);
198- 
199- ret = aclnnInnerQuantReduceScatterV2GetWorkspaceSize(mc2Context, x, scales, hcclBuffSize,
200- reduceOp, yDtype, worldSize, output,
201- workspaceSize, executor);
202- OP_LOGD("Execute QuantReduceScatterContext, aclnnGetWorkspaceSize ret is: %d.", ret);
203-#else
204- ret = aclnnInnerQuantReduceScatterGetWorkspaceSize(x, scales, group, reduceOp, yDtype,
205- worldSize, output, workspaceSize, executor);
206- OP_LOGD("Execute QuantReduceScatter, aclnnGetWorkspaceSize ret is: %d.", ret);
207-#endif
208- return ret;
209-}
210- 
211-extern "C" aclnnStatus aclnnQuantReduceScatterBase(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
212- const aclrtStream stream)
213-{
214- if (NnopbaseSetHcclServerType) {
215- NnopbaseSetHcclServerType(executor, NnopbaseHcclServerType::NNOPBASE_HCCL_SERVER_TYPE_MTE);
216- }
217- 
218-#if defined(BUILD_OPEN_PROJECT) && HCOMM_VERSION_NUM >= HCCL_CHANNEL_SUPPORT_VERSION
219- OP_LOGD("inner QuantReduceScatterContext start");
220- return aclnnInnerQuantReduceScatterV2(workspace, workspaceSize, executor, stream);
221-#else
222- OP_LOGD("inner QuantReduceScatter start");
223- return aclnnInnerQuantReduceScatter(workspace, workspaceSize, executor, stream);
224-#endif
225-}
Dmc2/quant_reduce_scatter/op_api/aclnn_quant_reduce_scatter_base.h+0-40
@@ -1,40 +0,0 @@
1-/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4- * CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9-**/
10- 
11-/*!
12- * \file aclnn_quant_reduce_scatter_base.h
13- * \brief
14- */
15-#ifndef OP_API_INC_QUANT_REDUCE_SCATTER_BASE_
16-#define OP_API_INC_QUANT_REDUCE_SCATTER_BASE_
17- 
18-#include <string>
19- 
20-#include "aclnn/aclnn_base.h"
21-#include "aclnn_util.h"
22-#include "hccl/hccl_types.h"
23- 
24-#ifdef __cplusplus
25-extern "C" {
26-#endif
27- 
28-ACLNN_API aclnnStatus aclnnQuantReduceScatterBaseGetWorkspaceSize(const aclTensor* x, const aclTensor* scales,
29- const char* group, const char* reduceOp,
30- aclTensor* output, uint64_t* workspaceSize,
31- aclOpExecutor** executor);
32- 
33-ACLNN_API aclnnStatus aclnnQuantReduceScatterBase(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
34- const aclrtStream stream);
35- 
36-#ifdef __cplusplus
37-}
38-#endif
39- 
40-#endif // OP_API_INC_QUANT_REDUCE_SCATTER_BASE_
Mmc2/quant_reduce_scatter/op_host/op_tiling/common/quant_reduce_scatter_util_tiling.cpp+76-107
@@ -34,21 +34,18 @@ static bool IsContains(const std::vector<uint32_t> &list, uint32_t value)
34 * @param context: 框架根据input，output，attrs等信息生成tiling需要的context34 * @param context: 框架根据input，output，attrs等信息生成tiling需要的context
35 * @return35 * @return
36 */36 */
37-static ge::graphStatus CheckAttrsInfo(const gert::TilingContext *context, TilingRunInfo &runInfo,37+static ge::graphStatus CheckAttrsInfo(const gert::TilingContext *context, TilingRunInfo &runInfo)
38- const QuantReduceScatterConfig& config)
39{38{
40 const char *nodeName = context->GetNodeName();39 const char *nodeName = context->GetNodeName();
41 const gert::RuntimeAttrs *attrs = context->GetAttrs();40 const gert::RuntimeAttrs *attrs = context->GetAttrs();
42 OP_TILING_CHECK(attrs == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "attrs"), return ge::GRAPH_FAILED);41 OP_TILING_CHECK(attrs == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "attrs"), return ge::GRAPH_FAILED);
43- if (!config.isMc2Context) {42+ // 校验group是否为空
44- // 校验group是否为空43+ const char *groupPtr = attrs->GetAttrPointer<char>(GROUP_INDEX);
45- const char *groupPtr = attrs->GetAttrPointer<char>(GROUP_INDEX);44+ OP_TILING_CHECK(groupPtr == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "group"), return ge::GRAPH_FAILED);
46- OP_TILING_CHECK(groupPtr == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "group"), return ge::GRAPH_FAILED);45+ OP_TILING_CHECK(std::string(groupPtr).empty(), OP_LOGE_WITH_INVALID_INPUT(nodeName, "group"),
47- OP_TILING_CHECK(std::string(groupPtr).empty(), OP_LOGE_WITH_INVALID_INPUT(nodeName, "group"),46+ return ge::GRAPH_FAILED);
48- return ge::GRAPH_FAILED);47+ runInfo.groupPtr = groupPtr;
49- runInfo.groupPtr = groupPtr;48+ runInfo.group = std::string(groupPtr);
50- runInfo.group = std::string(groupPtr);
51- }
52 // 校验reduce_op的类型是否为sum49 // 校验reduce_op的类型是否为sum
53 const char *reduceOpPtr = attrs->GetAttrPointer<char>(REDUCE_OP_INDEX);50 const char *reduceOpPtr = attrs->GetAttrPointer<char>(REDUCE_OP_INDEX);
54 OP_TILING_CHECK(reduceOpPtr == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "reduce_op"), return ge::GRAPH_FAILED);51 OP_TILING_CHECK(reduceOpPtr == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "reduce_op"), return ge::GRAPH_FAILED);
@@ -72,27 +69,19 @@ static ge::graphStatus CheckAttrsInfo(const gert::TilingContext *context, Tiling
72 * @param runInfo: 封装的doTiling所需要的参数69 * @param runInfo: 封装的doTiling所需要的参数
73 * @return70 * @return
74 */71 */
75-static ge::graphStatus SetRankSize(const gert::TilingContext *context, TilingRunInfo &runInfo,72+static ge::graphStatus SetRankSize(const gert::TilingContext *context, TilingRunInfo &runInfo)
76- const QuantReduceScatterConfig& config)
77{73{
78 const char *nodeName = context->GetNodeName();74 const char *nodeName = context->GetNodeName();
79 // attrs在函数CheckAttrsInfo中已做校验75 // attrs在函数CheckAttrsInfo中已做校验
80 const gert::RuntimeAttrs *attrs = context->GetAttrs();76 const gert::RuntimeAttrs *attrs = context->GetAttrs();
81 const int64_t *rankSizePtr = attrs->GetAttrPointer<int64_t>(WORLD_SIZE_INDEX);77 const int64_t *rankSizePtr = attrs->GetAttrPointer<int64_t>(WORLD_SIZE_INDEX);
82- if (!config.isMc2Context) {78+ if (rankSizePtr == nullptr || *rankSizePtr == RANK_SIZE_DEFAULT) {
83- if (rankSizePtr == nullptr || *rankSizePtr == RANK_SIZE_DEFAULT) {79+ int64_t rankSize = 0;
84- int64_t rankSize = 0;80+ OP_TILING_CHECK(!mc2tiling::GetRankSize(nodeName, runInfo.groupPtr, rankSize),
85- OP_TILING_CHECK(!mc2tiling::GetRankSize(nodeName, runInfo.groupPtr, rankSize),81+ OP_LOGE(nodeName, "Get rankSize failed."),
86- OP_LOGE(nodeName, "Get rankSize failed."),
87- return ge::GRAPH_FAILED);
88- runInfo.rankSize = rankSize;
89- } else {
90- runInfo.rankSize = *rankSizePtr;
91- }
92- } else {
93- OP_TILING_CHECK(rankSizePtr == nullptr || *rankSizePtr == RANK_SIZE_DEFAULT,
94- OP_LOGE(nodeName, "The rankSize is null or invalid value."),
95 return ge::GRAPH_FAILED);82 return ge::GRAPH_FAILED);
83+ runInfo.rankSize = rankSize;
84+ } else {
96 runInfo.rankSize = *rankSizePtr;85 runInfo.rankSize = *rankSizePtr;
97 }86 }
98 OP_TILING_CHECK(std::find(RANK_SIZE_LIST.begin(), RANK_SIZE_LIST.end(), runInfo.rankSize) >= RANK_SIZE_LIST.end(),87 OP_TILING_CHECK(std::find(RANK_SIZE_LIST.begin(), RANK_SIZE_LIST.end(), runInfo.rankSize) >= RANK_SIZE_LIST.end(),
@@ -107,13 +96,12 @@ static ge::graphStatus SetRankSize(const gert::TilingContext *context, TilingRun
107 * @param runInfo: 封装的doTiling所需要的参数96 * @param runInfo: 封装的doTiling所需要的参数
108 * @return97 * @return
109 */98 */
110-static bool SetQuantMode(const gert::TilingContext *context, TilingRunInfo &runInfo,99+static bool SetQuantMode(const gert::TilingContext *context, TilingRunInfo &runInfo)
111- const QuantReduceScatterConfig& config)
112{100{
113 const char *nodeName = context->GetNodeName();101 const char *nodeName = context->GetNodeName();
114 // context->GetInputDesc在函数CheckTensorDataType中已经校验102 // context->GetInputDesc在函数CheckTensorDataType中已经校验
115- ge::DataType xDtype = context->GetInputDesc(config.X_INDEX)->GetDataType();103+ ge::DataType xDtype = context->GetInputDesc(X_INDEX)->GetDataType();
116- ge::DataType scalesDtype = context->GetInputDesc(config.SCALES_INDEX)->GetDataType();104+ ge::DataType scalesDtype = context->GetInputDesc(SCALES_INDEX)->GetDataType();
117 // 0: 无量化模式; 1: TG量化; 2: MX量化105 // 0: 无量化模式; 1: TG量化; 2: MX量化
118 uint32_t quantMode = 0;106 uint32_t quantMode = 0;
119 if (IsContains(X_DTYPE_LIST, xDtype) && scalesDtype == ge::DT_FLOAT) {107 if (IsContains(X_DTYPE_LIST, xDtype) && scalesDtype == ge::DT_FLOAT) {
@@ -137,21 +125,20 @@ static bool SetQuantMode(const gert::TilingContext *context, TilingRunInfo &runI
137 * @param runInfo: 封装的doTiling所需要的参数125 * @param runInfo: 封装的doTiling所需要的参数
138 * @return126 * @return
139 */127 */
140-static bool CheckTensorDataType(const gert::TilingContext *context, TilingRunInfo &runInfo,128+static bool CheckTensorDataType(const gert::TilingContext *context, TilingRunInfo &runInfo)
141- const QuantReduceScatterConfig& config)
142{129{
143 const char *nodeName = context->GetNodeName();130 const char *nodeName = context->GetNodeName();
144 // 校验x的dtype131 // 校验x的dtype
145- auto xDesc = context->GetInputDesc(config.X_INDEX);132+ auto xDesc = context->GetInputDesc(X_INDEX);
146 OP_TILING_CHECK(xDesc == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "x"), return false);133 OP_TILING_CHECK(xDesc == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "x"), return false);
147- ge::DataType xDtype = context->GetInputDesc(config.X_INDEX)->GetDataType();134+ ge::DataType xDtype = context->GetInputDesc(X_INDEX)->GetDataType();
148 OP_TILING_CHECK(!IsContains(X_DTYPE_LIST, xDtype),135 OP_TILING_CHECK(!IsContains(X_DTYPE_LIST, xDtype),
149 OP_LOGE_FOR_INVALID_DTYPE(nodeName, "x", Ops::Base::ToString(xDtype).c_str(), "int8/hifloat8/float8_e4m3fn/float8_e5m2"),136 OP_LOGE_FOR_INVALID_DTYPE(nodeName, "x", Ops::Base::ToString(xDtype).c_str(), "int8/hifloat8/float8_e4m3fn/float8_e5m2"),
150 return false);137 return false);
151 // 校验scales的dtype138 // 校验scales的dtype
152- auto scalesDesc = context->GetInputDesc(config.SCALES_INDEX);139+ auto scalesDesc = context->GetInputDesc(SCALES_INDEX);
153 OP_TILING_CHECK(scalesDesc == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "scales"), return false);140 OP_TILING_CHECK(scalesDesc == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "scales"), return false);
154- ge::DataType scalesDtype = context->GetInputDesc(config.SCALES_INDEX)->GetDataType();141+ ge::DataType scalesDtype = context->GetInputDesc(SCALES_INDEX)->GetDataType();
155 OP_TILING_CHECK(!IsContains(SCALES_DTYPE_LIST, scalesDtype),142 OP_TILING_CHECK(!IsContains(SCALES_DTYPE_LIST, scalesDtype),
156 OP_LOGE_FOR_INVALID_DTYPE(nodeName, "scales", Ops::Base::ToString(scalesDtype).c_str(), "float/float8_e8m0"),143 OP_LOGE_FOR_INVALID_DTYPE(nodeName, "scales", Ops::Base::ToString(scalesDtype).c_str(), "float/float8_e8m0"),
157 return false);144 return false);
@@ -163,7 +150,7 @@ static bool CheckTensorDataType(const gert::TilingContext *context, TilingRunInf
163 OP_LOGE_FOR_INVALID_DTYPE(nodeName, "output", Ops::Base::ToString(outputType).c_str(), "float16/bfloat16/float"),150 OP_LOGE_FOR_INVALID_DTYPE(nodeName, "output", Ops::Base::ToString(outputType).c_str(), "float16/bfloat16/float"),
164 return false);151 return false);
165 // 设置量化模式152 // 设置量化模式
166- OP_TILING_CHECK(!SetQuantMode(context, runInfo, config), OP_LOGE(nodeName, "get quantMode error."), return false);153+ OP_TILING_CHECK(!SetQuantMode(context, runInfo), OP_LOGE(nodeName, "get quantMode error."), return false);
167 return true;154 return true;
168}155}
169 156 
@@ -173,13 +160,12 @@ static bool CheckTensorDataType(const gert::TilingContext *context, TilingRunInf
173 * @param opType: 当前op类型160 * @param opType: 当前op类型
174 * @return161 * @return
175 */162 */
176-static bool CheckXDimValid(const gert::TilingContext *context, const OpType opType,163+static bool CheckXDimValid(const gert::TilingContext *context, const OpType opType)
177- const QuantReduceScatterConfig& config)
178{164{
179 (void)opType; // Reserved for future extension165 (void)opType; // Reserved for future extension
180 const char *nodeName = context->GetNodeName();166 const char *nodeName = context->GetNodeName();
181 // context->GetInputShape在函数CheckInputTensorDim中已经校验167 // context->GetInputShape在函数CheckInputTensorDim中已经校验
182- size_t xDimNum = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDimNum();168+ size_t xDimNum = context->GetInputShape(X_INDEX)->GetStorageShape().GetDimNum();
183 // quant_all_reduce和quant_reduce_scatter算子的x可能是2维或者3维，即x.shape(bs, h)或x.shape(b, s, h)169 // quant_all_reduce和quant_reduce_scatter算子的x可能是2维或者3维，即x.shape(bs, h)或x.shape(b, s, h)
184 bool inValidDimNum = (xDimNum != TWO_DIMS) && (xDimNum != THREE_DIMS);170 bool inValidDimNum = (xDimNum != TWO_DIMS) && (xDimNum != THREE_DIMS);
185 OP_TILING_CHECK(inValidDimNum,171 OP_TILING_CHECK(inValidDimNum,
@@ -196,13 +182,12 @@ static bool CheckXDimValid(const gert::TilingContext *context, const OpType opTy
196 * @param opType: 当前op类型182 * @param opType: 当前op类型
197 * @return183 * @return
198 */184 */
199-static bool CheckXShapeValid(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType,185+static bool CheckXShapeValid(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType)
200- const QuantReduceScatterConfig& config)
201{186{
202 (void)opType; // Reserved for future extension187 (void)opType; // Reserved for future extension
203 const char *nodeName = context->GetNodeName();188 const char *nodeName = context->GetNodeName();
204 // context->GetInputShape在函数CheckInputTensorDim中已经校验189 // context->GetInputShape在函数CheckInputTensorDim中已经校验
205- const gert::StorageShape *xShape = context->GetInputShape(config.X_INDEX);190+ const gert::StorageShape *xShape = context->GetInputShape(X_INDEX);
206 // 获取x各维度值191 // 获取x各维度值
207 size_t xDimNum = xShape->GetStorageShape().GetDimNum();192 size_t xDimNum = xShape->GetStorageShape().GetDimNum();
208 uint64_t xValueOne = xShape->GetStorageShape().GetDim(DIM_ZERO);193 uint64_t xValueOne = xShape->GetStorageShape().GetDim(DIM_ZERO);
@@ -215,7 +200,7 @@ static bool CheckXShapeValid(const gert::TilingContext *context, TilingRunInfo &
215 // 当x是3维时，x.shape = (B, S, H)200 // 当x是3维时，x.shape = (B, S, H)
216 if (xDimNum == THREE_DIMS) {201 if (xDimNum == THREE_DIMS) {
217 xValueBS = xValueOne * xValueTwo;202 xValueBS = xValueOne * xValueTwo;
218- xValueH = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_TWO);203+ xValueH = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_TWO);
219 emptyTensor = emptyTensor || xValueH == 0;204 emptyTensor = emptyTensor || xValueH == 0;
220 }205 }
221 206 
@@ -251,10 +236,9 @@ static bool CheckXShapeValid(const gert::TilingContext *context, TilingRunInfo &
251 * @return 计算出的正确scales维度向量236 * @return 计算出的正确scales维度向量
252 */237 */
253static std::vector<uint64_t> CalculateExpectedScalesShape(const gert::TilingContext *context, 238static std::vector<uint64_t> CalculateExpectedScalesShape(const gert::TilingContext *context,
254- TilingRunInfo &runInfo,239+ TilingRunInfo &runInfo)
255- const QuantReduceScatterConfig& config)
256{240{
257- const gert::StorageShape *xShape = context->GetInputShape(config.X_INDEX);241+ const gert::StorageShape *xShape = context->GetInputShape(X_INDEX);
258 242
259 // 获取x的维度和值243 // 获取x的维度和值
260 size_t xDimNum = xShape->GetStorageShape().GetDimNum();244 size_t xDimNum = xShape->GetStorageShape().GetDimNum();
@@ -321,10 +305,10 @@ static std::string FormatShape(const std::vector<uint64_t> &dims)
321 */305 */
322static bool CheckScalesValid(const gert::TilingContext *context, 306static bool CheckScalesValid(const gert::TilingContext *context,
323 const std::vector<uint64_t> &expectedScalesDims,307 const std::vector<uint64_t> &expectedScalesDims,
324- const TilingRunInfo &runInfo, const QuantReduceScatterConfig& config)308+ const TilingRunInfo &runInfo)
325{309{
326 const char *nodeName = context->GetNodeName();310 const char *nodeName = context->GetNodeName();
327- const gert::StorageShape *scalesShape = context->GetInputShape(config.SCALES_INDEX);311+ const gert::StorageShape *scalesShape = context->GetInputShape(SCALES_INDEX);
328 312
329 // 将quantMode转换为可读字符串313 // 将quantMode转换为可读字符串
330 const char* quantModeStr = "";314 const char* quantModeStr = "";
@@ -364,29 +348,26 @@ static bool CheckScalesValid(const gert::TilingContext *context,
364 * @param opType:当前op类型348 * @param opType:当前op类型
365 * @return349 * @return
366 */350 */
367-static bool CheckInputTensorDim(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType,351+static bool CheckInputTensorDim(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType)
368- const QuantReduceScatterConfig& config)
369{352{
370 const char *nodeName = context->GetNodeName();353 const char *nodeName = context->GetNodeName();
371 // 1.校验x相关354 // 1.校验x相关
372- const gert::StorageShape *xShape = context->GetInputShape(config.X_INDEX);355+ const gert::StorageShape *xShape = context->GetInputShape(X_INDEX);
373 // 校验x不为空356 // 校验x不为空
374 OP_TILING_CHECK(xShape == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "xShape"), return false);357 OP_TILING_CHECK(xShape == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "xShape"), return false);
375 // 校验x维度数量合法性358 // 校验x维度数量合法性
376- OP_TILING_CHECK(!CheckXDimValid(context, opType, config),359+ OP_TILING_CHECK(!CheckXDimValid(context, opType), OP_LOGE(nodeName, "x dimensions is invalid."), return false);
377- OP_LOGE(nodeName, "x dimensions is invalid."), return false);
378 // 校验x.shape合法性360 // 校验x.shape合法性
379- OP_TILING_CHECK(!CheckXShapeValid(context, runInfo, opType, config),361+ OP_TILING_CHECK(!CheckXShapeValid(context, runInfo, opType), OP_LOGE(nodeName, "x shapes is invalid."), return false);
380- OP_LOGE(nodeName, "x shapes is invalid."), return false);
381 362 
382 // 2.校验scales363 // 2.校验scales
383- const gert::StorageShape *scalesShape = context->GetInputShape(config.SCALES_INDEX);364+ const gert::StorageShape *scalesShape = context->GetInputShape(SCALES_INDEX);
384 // 根据x计算正确的scales, 当scale形状不匹配时，会打印预期的形状和实际的形状365 // 根据x计算正确的scales, 当scale形状不匹配时，会打印预期的形状和实际的形状
385- std::vector<uint64_t> expectedScalesDims = CalculateExpectedScalesShape(context, runInfo, config);366+ std::vector<uint64_t> expectedScalesDims = CalculateExpectedScalesShape(context, runInfo);
386 // 校验scales不为空367 // 校验scales不为空
387 OP_TILING_CHECK(scalesShape == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "scales"), return false);368 OP_TILING_CHECK(scalesShape == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "scales"), return false);
388 // 校验scales维度和shape是否正确369 // 校验scales维度和shape是否正确
389- OP_TILING_CHECK(!CheckScalesValid(context, expectedScalesDims, runInfo, config),370+ OP_TILING_CHECK(!CheckScalesValid(context, expectedScalesDims, runInfo),
390 OP_LOGE(nodeName, "scales dimensions and shapes is invalid in the quantmode."), return false);371 OP_LOGE(nodeName, "scales dimensions and shapes is invalid in the quantmode."), return false);
391 return true;372 return true;
392}373}
@@ -419,14 +400,13 @@ static bool CheckOutputDimSize(const gert::TilingContext *context, size_t output
419 * @brief 检查quant_all_reduce的输出形状400 * @brief 检查quant_all_reduce的输出形状
420 */401 */
421static bool CheckAllReduceOutputShape(const gert::TilingContext *context, const gert::StorageShape *outputShape,402static bool CheckAllReduceOutputShape(const gert::TilingContext *context, const gert::StorageShape *outputShape,
422- size_t outputDim, size_t xDimNum, TilingRunInfo &runInfo, const char *nodeName,403+ size_t outputDim, size_t xDimNum, TilingRunInfo &runInfo, const char *nodeName)
423- const QuantReduceScatterConfig& config)
424{404{
425 (void)xDimNum; // Reserved for future extension405 (void)xDimNum; // Reserved for future extension
426 uint64_t outputValueOne = outputShape->GetStorageShape().GetDim(DIM_ZERO);406 uint64_t outputValueOne = outputShape->GetStorageShape().GetDim(DIM_ZERO);
427 uint64_t outputValueTwo = outputShape->GetStorageShape().GetDim(DIM_ONE);407 uint64_t outputValueTwo = outputShape->GetStorageShape().GetDim(DIM_ONE);
428- uint64_t xValueOne = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);408+ uint64_t xValueOne = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);
429- uint64_t xValueTwo = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);409+ uint64_t xValueTwo = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
430 410
431 // 对于quant_all_reduce算子，output.shape必须等于x.shape411 // 对于quant_all_reduce算子，output.shape必须等于x.shape
432 bool invalidShape = (xValueOne != outputValueOne) || (xValueTwo != outputValueTwo); // 校验前两维的大小412 bool invalidShape = (xValueOne != outputValueOne) || (xValueTwo != outputValueTwo); // 校验前两维的大小
@@ -434,7 +414,7 @@ static bool CheckAllReduceOutputShape(const gert::TilingContext *context, const
434 // quant_all_reduce算子支持三维，output可能需要校验第3维414 // quant_all_reduce算子支持三维，output可能需要校验第3维
435 if (outputDim == THREE_DIMS) {415 if (outputDim == THREE_DIMS) {
436 uint64_t outputValueThree = outputShape->GetStorageShape().GetDim(DIM_TWO);416 uint64_t outputValueThree = outputShape->GetStorageShape().GetDim(DIM_TWO);
437- uint64_t xValueThree = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_TWO);417+ uint64_t xValueThree = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_TWO);
438 OP_LOGI(nodeName, "output dim2 is %lu, x dim2 is %lu", outputValueThree, xValueThree);418 OP_LOGI(nodeName, "output dim2 is %lu, x dim2 is %lu", outputValueThree, xValueThree);
439 invalidShape = invalidShape || (xValueThree != outputValueThree); // 校验第三维的大小419 invalidShape = invalidShape || (xValueThree != outputValueThree); // 校验第三维的大小
440 OP_TILING_CHECK(invalidShape,420 OP_TILING_CHECK(invalidShape,
@@ -459,12 +439,11 @@ static bool CheckAllReduceOutputShape(const gert::TilingContext *context, const
459static bool CheckReduceScatter3DShape(const gert::TilingContext *context,439static bool CheckReduceScatter3DShape(const gert::TilingContext *context,
460 uint64_t outputValueOne, uint64_t outputValueTwo,440 uint64_t outputValueOne, uint64_t outputValueTwo,
461 uint64_t xValueOne, uint64_t xValueTwo,441 uint64_t xValueOne, uint64_t xValueTwo,
462- TilingRunInfo &runInfo, const char *nodeName,442+ TilingRunInfo &runInfo, const char *nodeName)
463- const QuantReduceScatterConfig& config)
464{443{
465 // 若X为3维，则要对b,s进行合轴，再与output判断是否合法444 // 若X为3维，则要对b,s进行合轴，再与output判断是否合法
466 uint64_t xValueBS = xValueOne * xValueTwo;445 uint64_t xValueBS = xValueOne * xValueTwo;
467- uint64_t xValueThree = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_TWO);446+ uint64_t xValueThree = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_TWO);
468 bool invalidShape = xValueBS / runInfo.rankSize != outputValueOne; // 校验bs轴447 bool invalidShape = xValueBS / runInfo.rankSize != outputValueOne; // 校验bs轴
469 invalidShape = invalidShape || (xValueThree != outputValueTwo); // 校验h轴448 invalidShape = invalidShape || (xValueThree != outputValueTwo); // 校验h轴
470 OP_TILING_CHECK(invalidShape,449 OP_TILING_CHECK(invalidShape,
@@ -497,19 +476,18 @@ static bool CheckReduceScatter2DShape(uint64_t outputValueOne, uint64_t outputVa
497 * @brief 检查quant_reduce_scatter的输出形状476 * @brief 检查quant_reduce_scatter的输出形状
498 */477 */
499static bool CheckReduceScatterOutputShape(const gert::TilingContext *context, const gert::StorageShape *outputShape,478static bool CheckReduceScatterOutputShape(const gert::TilingContext *context, const gert::StorageShape *outputShape,
500- size_t outputDim, size_t xDimNum, TilingRunInfo &runInfo,479+ size_t outputDim, size_t xDimNum, TilingRunInfo &runInfo, const char *nodeName)
501- const char *nodeName, const QuantReduceScatterConfig& config)
502{480{
503 (void)outputDim; // Reserved for future extension481 (void)outputDim; // Reserved for future extension
504 uint64_t outputValueOne = outputShape->GetStorageShape().GetDim(DIM_ZERO);482 uint64_t outputValueOne = outputShape->GetStorageShape().GetDim(DIM_ZERO);
505 uint64_t outputValueTwo = outputShape->GetStorageShape().GetDim(DIM_ONE);483 uint64_t outputValueTwo = outputShape->GetStorageShape().GetDim(DIM_ONE);
506- uint64_t xValueOne = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);484+ uint64_t xValueOne = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);
507- uint64_t xValueTwo = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);485+ uint64_t xValueTwo = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
508 486
509 // 对于quant_reduce_scatter算子, 输出output一定为2维，判断x维度大小决定是否b,s合轴487 // 对于quant_reduce_scatter算子, 输出output一定为2维，判断x维度大小决定是否b,s合轴
510 if (xDimNum == THREE_DIMS) {488 if (xDimNum == THREE_DIMS) {
511 return CheckReduceScatter3DShape(context, outputValueOne, outputValueTwo, 489 return CheckReduceScatter3DShape(context, outputValueOne, outputValueTwo,
512- xValueOne, xValueTwo, runInfo, nodeName, config);490+ xValueOne, xValueTwo, runInfo, nodeName);
513 } else {491 } else {
514 return CheckReduceScatter2DShape(outputValueOne, outputValueTwo, 492 return CheckReduceScatter2DShape(outputValueOne, outputValueTwo,
515 xValueOne, xValueTwo, runInfo, nodeName);493 xValueOne, xValueTwo, runInfo, nodeName);
@@ -522,15 +500,14 @@ static bool CheckReduceScatterOutputShape(const gert::TilingContext *context, co
522 * @param runInfo: 封装的doTiling所需要的参数500 * @param runInfo: 封装的doTiling所需要的参数
523 * @return501 * @return
524 */502 */
525-static bool CheckOutputDim(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType,503+static bool CheckOutputDim(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType)
526- const QuantReduceScatterConfig& config)
527{504{
528 const char *nodeName = context->GetNodeName();505 const char *nodeName = context->GetNodeName();
529 // context->GetOutputShape在函数CheckOutputTensorDim中已经校验506 // context->GetOutputShape在函数CheckOutputTensorDim中已经校验
530 const gert::StorageShape *outputShape = context->GetOutputShape(OUTPUT_INDEX);507 const gert::StorageShape *outputShape = context->GetOutputShape(OUTPUT_INDEX);
531 size_t outputDim = outputShape->GetStorageShape().GetDimNum();508 size_t outputDim = outputShape->GetStorageShape().GetDimNum();
532 // context->GetInputShape在函数CheckInputTensorDim中已经校验509 // context->GetInputShape在函数CheckInputTensorDim中已经校验
533- size_t xDimNum = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDimNum();510+ size_t xDimNum = context->GetInputShape(X_INDEX)->GetStorageShape().GetDimNum();
534 // 检查output的维度大小 511 // 检查output的维度大小
535 if (!CheckOutputDimSize(context, outputDim, xDimNum, opType, nodeName)) {512 if (!CheckOutputDimSize(context, outputDim, xDimNum, opType, nodeName)) {
536 return false;513 return false;
@@ -538,9 +515,9 @@ static bool CheckOutputDim(const gert::TilingContext *context, TilingRunInfo &ru
538 515 
539 // 检查输出output形状与输入x形状的关系516 // 检查输出output形状与输入x形状的关系
540 if (opType == OpType::OP_QUANT_ALL_REDUCE) {517 if (opType == OpType::OP_QUANT_ALL_REDUCE) {
541- return CheckAllReduceOutputShape(context, outputShape, outputDim, xDimNum, runInfo, nodeName, config);518+ return CheckAllReduceOutputShape(context, outputShape, outputDim, xDimNum, runInfo, nodeName);
542 } else {519 } else {
543- return CheckReduceScatterOutputShape(context, outputShape, outputDim, xDimNum, runInfo, nodeName, config);520+ return CheckReduceScatterOutputShape(context, outputShape, outputDim, xDimNum, runInfo, nodeName);
544 }521 }
545}522}
546 523 
@@ -551,14 +528,13 @@ static bool CheckOutputDim(const gert::TilingContext *context, TilingRunInfo &ru
551 * @param opType: 当前op类型528 * @param opType: 当前op类型
552 * @return529 * @return
553 */530 */
554-static bool CheckOutputTensorDim(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType,531+static bool CheckOutputTensorDim(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType)
555- const QuantReduceScatterConfig& config)
556{532{
557 const char *nodeName = context->GetNodeName();533 const char *nodeName = context->GetNodeName();
558 // 红线校验534 // 红线校验
559 const gert::StorageShape *outputShape = context->GetOutputShape(OUTPUT_INDEX);535 const gert::StorageShape *outputShape = context->GetOutputShape(OUTPUT_INDEX);
560 OP_TILING_CHECK(outputShape == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "output"), return false);536 OP_TILING_CHECK(outputShape == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "output"), return false);
561- return CheckOutputDim(context, runInfo, opType, config);537+ return CheckOutputDim(context, runInfo, opType);
562}538}
563 539 
564/**540/**
@@ -566,17 +542,17 @@ static bool CheckOutputTensorDim(const gert::TilingContext *context, TilingRunIn
566 * @param context: 框架根据input，output，attrs等信息生成tiling需要的context542 * @param context: 框架根据input，output，attrs等信息生成tiling需要的context
567 * @return543 * @return
568 */544 */
569-static bool CheckTensorFormat(const gert::TilingContext *context, const QuantReduceScatterConfig& config)545+static bool CheckTensorFormat(const gert::TilingContext *context)
570{546{
571 const char *nodeName = context->GetNodeName();547 const char *nodeName = context->GetNodeName();
572 // context->GetInputDesc在CheckTensorDataType函数中已经校验548 // context->GetInputDesc在CheckTensorDataType函数中已经校验
573- auto xDesc = context->GetInputDesc(config.X_INDEX);549+ auto xDesc = context->GetInputDesc(X_INDEX);
574 ge::Format xFormat = static_cast<ge::Format>(ge::GetPrimaryFormat(xDesc->GetStorageFormat()));550 ge::Format xFormat = static_cast<ge::Format>(ge::GetPrimaryFormat(xDesc->GetStorageFormat()));
575 OP_TILING_CHECK(551 OP_TILING_CHECK(
576 xFormat != ge::FORMAT_ND,552 xFormat != ge::FORMAT_ND,
577 OP_LOGE_FOR_INVALID_FORMAT(nodeName, "x", Ops::Base::ToString(xFormat).c_str(), "ND"),553 OP_LOGE_FOR_INVALID_FORMAT(nodeName, "x", Ops::Base::ToString(xFormat).c_str(), "ND"),
578 return false);554 return false);
579- auto scalesDesc = context->GetInputDesc(config.SCALES_INDEX);555+ auto scalesDesc = context->GetInputDesc(SCALES_INDEX);
580 ge::Format scalesFormat = static_cast<ge::Format>(ge::GetPrimaryFormat(scalesDesc->GetStorageFormat()));556 ge::Format scalesFormat = static_cast<ge::Format>(ge::GetPrimaryFormat(scalesDesc->GetStorageFormat()));
581 OP_TILING_CHECK(scalesFormat != ge::FORMAT_ND,557 OP_TILING_CHECK(scalesFormat != ge::FORMAT_ND,
582 OP_LOGE_FOR_INVALID_FORMAT(nodeName, "scales", Ops::Base::ToString(scalesFormat).c_str(), "ND"),558 OP_LOGE_FOR_INVALID_FORMAT(nodeName, "scales", Ops::Base::ToString(scalesFormat).c_str(), "ND"),
@@ -596,25 +572,24 @@ static bool CheckTensorFormat(const gert::TilingContext *context, const QuantRed
596 * @param runInfo: 封装的doTiling所需要的参数572 * @param runInfo: 封装的doTiling所需要的参数
597 * @return573 * @return
598 */574 */
599-static bool CheckWindowSize(const gert::TilingContext *context, const TilingRunInfo &runInfo,575+static bool CheckWindowSize(const gert::TilingContext *context, const TilingRunInfo &runInfo)
600- const QuantReduceScatterConfig& config)
601{576{
602 const char *nodeName = context->GetNodeName();577 const char *nodeName = context->GetNodeName();
603 // 获取量化模式，数据类型578 // 获取量化模式，数据类型
604- uint64_t xValueOne = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);579+ uint64_t xValueOne = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);
605- uint64_t xValueTwo = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);580+ uint64_t xValueTwo = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
606- uint64_t scalesValueOne = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_ZERO);581+ uint64_t scalesValueOne = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(DIM_ZERO);
607- uint64_t scalesValueTwo = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_ONE);582+ uint64_t scalesValueTwo = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(DIM_ONE);
608 583 
609 // 计算xDataSize584 // 计算xDataSize
610 uint64_t xValue = xValueOne * xValueTwo;585 uint64_t xValue = xValueOne * xValueTwo;
611 uint64_t scalesValue = scalesValueOne * scalesValueTwo;586 uint64_t scalesValue = scalesValueOne * scalesValueTwo;
612 uint32_t scalesLastDim = DIM_TWO;587 uint32_t scalesLastDim = DIM_TWO;
613- size_t xDimNum = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDimNum();588+ size_t xDimNum = context->GetInputShape(X_INDEX)->GetStorageShape().GetDimNum();
614 if (xDimNum == THREE_DIMS) {589 if (xDimNum == THREE_DIMS) {
615- uint64_t xValueThree = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_TWO);590+ uint64_t xValueThree = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_TWO);
616 xValue = xValue * xValueThree;591 xValue = xValue * xValueThree;
617- uint64_t scalesValueThree = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_TWO);592+ uint64_t scalesValueThree = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(DIM_TWO);
618 scalesValue = scalesValue * scalesValueThree;593 scalesValue = scalesValue * scalesValueThree;
619 scalesLastDim = DIM_THREE;594 scalesLastDim = DIM_THREE;
620 }595 }
@@ -628,7 +603,7 @@ static bool CheckWindowSize(const gert::TilingContext *context, const TilingRunI
628 scalesSize = scalesValue * SCALE_DTYPE_SIZE_FOUR;603 scalesSize = scalesValue * SCALE_DTYPE_SIZE_FOUR;
629 } else if (runInfo.quantMode == MX_QUANT_MOD) {604 } else if (runInfo.quantMode == MX_QUANT_MOD) {
630 // scales的最后一维一定为2605 // scales的最后一维一定为2
631- uint64_t scalesValueLast = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(scalesLastDim);606+ uint64_t scalesValueLast = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(scalesLastDim);
632 scalesSize = scalesValue * scalesValueLast * SCALE_DTYPE_SIZE_ONE;607 scalesSize = scalesValue * scalesValueLast * SCALE_DTYPE_SIZE_ONE;
633 }608 }
634 uint64_t scalesDataSize = ((scalesSize + WIN_ADDR_ALIGN - 1UL) / WIN_ADDR_ALIGN) * WIN_ADDR_ALIGN;609 uint64_t scalesDataSize = ((scalesSize + WIN_ADDR_ALIGN - 1UL) / WIN_ADDR_ALIGN) * WIN_ADDR_ALIGN;
@@ -688,31 +663,25 @@ ge::graphStatus QuantReduceScatterUtilTiling::CheckNpuArch(const gert::TilingCon
688 * @return663 * @return
689 */664 */
690ge::graphStatus QuantReduceScatterUtilTiling::CheckTilingFunc(gert::TilingContext *context, TilingRunInfo &runInfo,665ge::graphStatus QuantReduceScatterUtilTiling::CheckTilingFunc(gert::TilingContext *context, TilingRunInfo &runInfo,
691- const OpType opType,666+ const OpType opType)
692- const QuantReduceScatterConfig& config)
693{667{
694 const char *nodeName = context->GetNodeName();668 const char *nodeName = context->GetNodeName();
695 // set group669 // set group
696- OP_TILING_CHECK(CheckAttrsInfo(context, runInfo, config) != ge::GRAPH_SUCCESS,670+ OP_TILING_CHECK(CheckAttrsInfo(context, runInfo) != ge::GRAPH_SUCCESS, OP_LOGE(nodeName, "attrs are invalied."),
697- OP_LOGE(nodeName, "attrs are invalied."),
698 return ge::GRAPH_FAILED);671 return ge::GRAPH_FAILED);
699 // set rankSize672 // set rankSize
700- OP_TILING_CHECK(SetRankSize(context, runInfo, config) != ge::GRAPH_SUCCESS,673+ OP_TILING_CHECK(SetRankSize(context, runInfo) != ge::GRAPH_SUCCESS, OP_LOGE(nodeName, "set rankSize failed."),
701- OP_LOGE(nodeName, "set rankSize failed."),
702 return ge::GRAPH_FAILED);674 return ge::GRAPH_FAILED);
703 // set quantMode675 // set quantMode
704- OP_TILING_CHECK(!CheckTensorDataType(context, runInfo, config),676+ OP_TILING_CHECK(!CheckTensorDataType(context, runInfo), OP_LOGE(nodeName, "tensor datatype is invalid."),
705- OP_LOGE(nodeName, "tensor datatype is invalid."),
706 return ge::GRAPH_FAILED);677 return ge::GRAPH_FAILED);
707- OP_TILING_CHECK(!CheckInputTensorDim(context, runInfo, opType, config),678+ OP_TILING_CHECK(!CheckInputTensorDim(context, runInfo, opType), OP_LOGE(nodeName, "input tensor dim is invalid."),
708- OP_LOGE(nodeName, "input tensor dim is invalid."),
709 return ge::GRAPH_FAILED);679 return ge::GRAPH_FAILED);
710- OP_TILING_CHECK(!CheckOutputTensorDim(context, runInfo, opType, config),680+ OP_TILING_CHECK(!CheckOutputTensorDim(context, runInfo, opType), OP_LOGE(nodeName, "output tensor dim is invalid."),
711- OP_LOGE(nodeName, "output tensor dim is invalid."),
712 return ge::GRAPH_FAILED);681 return ge::GRAPH_FAILED);
713- OP_TILING_CHECK(!CheckTensorFormat(context, config), OP_LOGE(nodeName, "tensor format is invalid."),682+ OP_TILING_CHECK(!CheckTensorFormat(context), OP_LOGE(nodeName, "tensor format is invalid."),
714 return ge::GRAPH_FAILED);683 return ge::GRAPH_FAILED);
715- OP_TILING_CHECK(!CheckWindowSize(context, runInfo, config), OP_LOGE(nodeName, "HCCL_BUFFSIZE is too small."),684+ OP_TILING_CHECK(!CheckWindowSize(context, runInfo), OP_LOGE(nodeName, "HCCL_BUFFSIZE is too small."),
716 return ge::GRAPH_FAILED);685 return ge::GRAPH_FAILED);
717 OP_TILING_CHECK(SetWorkSpace(context) != ge::GRAPH_SUCCESS, OP_LOGE(nodeName, "set workspace failed."),686 OP_TILING_CHECK(SetWorkSpace(context) != ge::GRAPH_SUCCESS, OP_LOGE(nodeName, "set workspace failed."),
718 return ge::GRAPH_FAILED);687 return ge::GRAPH_FAILED);
Mmc2/quant_reduce_scatter/op_host/op_tiling/common/quant_reduce_scatter_util_tiling.h+3-8
@@ -21,12 +21,8 @@ using namespace ge;
21using namespace gert;21using namespace gert;
22 22 
23// input index23// input index
24-struct QuantReduceScatterConfig {24+constexpr size_t X_INDEX = 0;
25- uint64_t CONTEXT_INDEX = 0;25+constexpr size_t SCALES_INDEX = 1;
26- uint64_t X_INDEX = 0;
27- uint64_t SCALES_INDEX = 1;
28- bool isMc2Context = false;
29-};
30// output index26// output index
31constexpr size_t OUTPUT_INDEX = 0;27constexpr size_t OUTPUT_INDEX = 0;
32// attr index28// attr index
@@ -88,8 +84,7 @@ struct TilingRunInfo {
88class QuantReduceScatterUtilTiling {84class QuantReduceScatterUtilTiling {
89public:85public:
90 static ge::graphStatus CheckNpuArch(const gert::TilingContext *context);86 static ge::graphStatus CheckNpuArch(const gert::TilingContext *context);
91- static ge::graphStatus CheckTilingFunc(gert::TilingContext *context, TilingRunInfo &runInfo,87+ static ge::graphStatus CheckTilingFunc(gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType);
92- const OpType opType, const QuantReduceScatterConfig& config);
93};88};
94 89 
95}; // namespace MC2Tiling90}; // namespace MC2Tiling
Mmc2/quant_reduce_scatter/op_host/op_tiling/quant_reduce_scatter_tiling.cpp+11-18
@@ -67,8 +67,7 @@ static ge::graphStatus SetHcommCfg(const gert::TilingContext *context, QuantRedu
67 * @param tilingData: 框架根据context的opName匹配tiling模板，计算产生的tilingData67 * @param tilingData: 框架根据context的opName匹配tiling模板，计算产生的tilingData
68 * @return68 * @return
69 */69 */
70-static void SetTilingData(gert::TilingContext *context, QuantReduceScatterTilingData &tilingData,70+static void SetTilingData(gert::TilingContext *context, QuantReduceScatterTilingData &tilingData)
71- const QuantReduceScatterConfig& config)
72{71{
73 fe::PlatFormInfos *platformInfoPtr = context->GetPlatformInfo();72 fe::PlatFormInfos *platformInfoPtr = context->GetPlatformInfo();
74 platform_ascendc::PlatformAscendC ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);73 platform_ascendc::PlatformAscendC ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
@@ -76,20 +75,19 @@ static void SetTilingData(gert::TilingContext *context, QuantReduceScatterTiling
76 uint32_t aivNum = ascendcPlatform.GetCoreNumAiv();75 uint32_t aivNum = ascendcPlatform.GetCoreNumAiv();
77 context->SetBlockDim(ascendcPlatform.CalcTschBlockDim(aivNum, 0, aivNum));76 context->SetBlockDim(ascendcPlatform.CalcTschBlockDim(aivNum, 0, aivNum));
78 tilingData.quantReduceScatterTilingInfo.aivNum = aivNum;77 tilingData.quantReduceScatterTilingInfo.aivNum = aivNum;
79- uint64_t xValueBS = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);78+ uint64_t xValueBS = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);
80- uint64_t xValueH = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);79+ uint64_t xValueH = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
81- uint64_t scalesValueH = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_ONE);80+ uint64_t scalesValueH = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(DIM_ONE);
82 // 3d场景，context->GetInputShape在函数CheckInputTensorDim中已经校验81 // 3d场景，context->GetInputShape在函数CheckInputTensorDim中已经校验
83- if (context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDimNum() == THREE_DIMS) {82+ if (context->GetInputShape(X_INDEX)->GetStorageShape().GetDimNum() == THREE_DIMS) {
84- xValueBS = xValueBS * context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);83+ xValueBS = xValueBS * context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
85- xValueH = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_TWO);84+ xValueH = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_TWO);
86- scalesValueH = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_TWO);85+ scalesValueH = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(DIM_TWO);
87 }86 }
88 tilingData.quantReduceScatterTilingInfo.bs = xValueBS;87 tilingData.quantReduceScatterTilingInfo.bs = xValueBS;
89 tilingData.quantReduceScatterTilingInfo.hiddenSize = xValueH;88 tilingData.quantReduceScatterTilingInfo.hiddenSize = xValueH;
90 tilingData.quantReduceScatterTilingInfo.scaleHiddenSize = scalesValueH;89 tilingData.quantReduceScatterTilingInfo.scaleHiddenSize = scalesValueH;
91 tilingData.quantReduceScatterTilingInfo.totalWinSize = mc2tiling::Mc2TilingUtils::GetMaxWindowSize();90 tilingData.quantReduceScatterTilingInfo.totalWinSize = mc2tiling::Mc2TilingUtils::GetMaxWindowSize();
92- tilingData.quantReduceScatterTilingInfo.isMc2Context = config.isMc2Context;
93}91}
94 92 
95// 基于 TARGET_ITER 公式计算 host 推荐的 xPerBlock（先除 rankSize 再反推），写入 tilingData93// 基于 TARGET_ITER 公式计算 host 推荐的 xPerBlock（先除 rankSize 再反推），写入 tilingData
@@ -148,18 +146,13 @@ static ge::graphStatus QuantReduceScatterTilingFunc(gert::TilingContext *context
148 OP_TILING_CHECK(QuantReduceScatterUtilTiling::CheckNpuArch(context) != ge::GRAPH_SUCCESS,146 OP_TILING_CHECK(QuantReduceScatterUtilTiling::CheckNpuArch(context) != ge::GRAPH_SUCCESS,
149 OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(nodeName, "npuArch", "non-DAV_3510", "The value of npuArch must be DAV_3510"),147 OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(nodeName, "npuArch", "non-DAV_3510", "The value of npuArch must be DAV_3510"),
150 return ge::GRAPH_FAILED);148 return ge::GRAPH_FAILED);
151- 149+ OP_TILING_CHECK(QuantReduceScatterUtilTiling::CheckTilingFunc(context, runInfo, OpType::OP_QUANT_REDUCE_SCATTER) !=
152- QuantReduceScatterConfig config;150+ ge::GRAPH_SUCCESS,
153- config.X_INDEX = 0;
154- config.SCALES_INDEX = 1;
155- config.isMc2Context = false;
156- OP_TILING_CHECK(QuantReduceScatterUtilTiling::CheckTilingFunc(context, runInfo,
157- OpType::OP_QUANT_REDUCE_SCATTER, config) != ge::GRAPH_SUCCESS,
158 OP_LOGE(nodeName, "tiling check failed in quant_reduce_scatter."), return ge::GRAPH_FAILED);151 OP_LOGE(nodeName, "tiling check failed in quant_reduce_scatter."), return ge::GRAPH_FAILED);
159 152 
160 OP_TILING_CHECK(SetHcommCfg(context, tilingData, runInfo) != ge::GRAPH_SUCCESS,153 OP_TILING_CHECK(SetHcommCfg(context, tilingData, runInfo) != ge::GRAPH_SUCCESS,
161 OP_LOGE(nodeName, "SetHCommCfg failed."), return ge::GRAPH_FAILED);154 OP_LOGE(nodeName, "SetHCommCfg failed."), return ge::GRAPH_FAILED);
162- SetTilingData(context, *tilingData, config);155+ SetTilingData(context, *tilingData);
163 SetXPerBlock(*tilingData, runInfo);156 SetXPerBlock(*tilingData, runInfo);
164 SetTilingKey(context);157 SetTilingKey(context);
165 PrintTilingDataInfo(context, *tilingData);158 PrintTilingDataInfo(context, *tilingData);
Mmc2/quant_reduce_scatter/op_kernel/mte_comm.h+16-54
@@ -18,15 +18,11 @@
18 18 
19#include "adv_api/hccl/hccl.h"19#include "adv_api/hccl/hccl.h"
20#include "adv_api/reduce/sum.h"20#include "adv_api/reduce/sum.h"
21-#include "quant_reduce_scatter_tiling_data.h"
22#include "../../common/op_kernel/moe_distribute_base.h"21#include "../../common/op_kernel/moe_distribute_base.h"
23-#include "../../common/op_kernel/mc2_quant_reduce_scatter_context.h"
24 22 
25namespace QuantMTECommImpl {23namespace QuantMTECommImpl {
26 24 
27using namespace AscendC;25using namespace AscendC;
28-using namespace Mc2Aclnn;
29- 
30// 后缀_BYTES 表示单位为字节大小B, _NUM 表示单位为个26// 后缀_BYTES 表示单位为字节大小B, _NUM 表示单位为个
31constexpr static uint32_t UB_ALIGN_BYTES = 32U; // UB按32B对齐27constexpr static uint32_t UB_ALIGN_BYTES = 32U; // UB按32B对齐
32constexpr uint32_t FLOAT_UB_ALIGN_NUM = 8U; // float格式下32B对齐需要 32/4 =8个28constexpr uint32_t FLOAT_UB_ALIGN_NUM = 8U; // float格式下32B对齐需要 32/4 =8个
@@ -45,7 +41,6 @@ template<TemplateTypeClass>
45class MTECommunication {41class MTECommunication {
46public:42public:
47 __aicore__ inline MTECommunication() {};43 __aicore__ inline MTECommunication() {};
48- __aicore__ inline void InitMc2Context(GM_ADDR mc2Context, const QuantReduceScatterTilingData *tilingData);
49 __aicore__ inline void InitHcclContext();44 __aicore__ inline void InitHcclContext();
50 __aicore__ inline void InitParams();45 __aicore__ inline void InitParams();
51 __aicore__ inline void InitGMTensor(GM_ADDR x, GM_ADDR scales, GM_ADDR output, uint64_t alignedXSize, uint64_t dataSpaceGmSize);46 __aicore__ inline void InitGMTensor(GM_ADDR x, GM_ADDR scales, GM_ADDR output, uint64_t alignedXSize, uint64_t dataSpaceGmSize);
@@ -62,12 +57,7 @@ public:
62 __aicore__ inline GM_ADDR GetWinDataAddrGm(uint32_t rankId, uint32_t winFlag);57 __aicore__ inline GM_ADDR GetWinDataAddrGm(uint32_t rankId, uint32_t winFlag);
63 __aicore__ inline GM_ADDR GetWinStatusAddrGm(uint32_t rankId, uint32_t winFlag);58 __aicore__ inline GM_ADDR GetWinStatusAddrGm(uint32_t rankId, uint32_t winFlag);
64 59 
65- uint32_t rankIdHccl_{0};60+ __gm__ Mc2Kernel::HcclOpParam *hcclContext_;
66- uint32_t rankDimHccl_{0};
67- bool isMc2Context_ = false;
68- __gm__ Mc2QuantReduceScatterContext* mc2Context_{nullptr};
69- __gm__ Mc2Kernel::HcclOpParam *hcclContext_{nullptr};
70- 
71 uint32_t aivId_{0};61 uint32_t aivId_{0};
72 uint64_t aivNum_{0};62 uint64_t aivNum_{0};
73 uint32_t round_{0};63 uint32_t round_{0};
@@ -105,29 +95,10 @@ private:
105 TBuf<> stateResetBuf_;95 TBuf<> stateResetBuf_;
106};96};
107 97 
108-template <TemplateTypeClass>
109-__aicore__ inline void MTECommunication<TemplateType>::InitMc2Context(
110- GM_ADDR mc2Context, const QuantReduceScatterTilingData *tilingData)
111-{
112- isMc2Context_ = tilingData->quantReduceScatterTilingInfo.isMc2Context;
113- if (isMc2Context_) {
114- mc2Context_ = (__gm__ Mc2QuantReduceScatterContext*)mc2Context;
115- rankIdHccl_ = mc2Context_->rankId;
116- rankDimHccl_ = mc2Context_->rankDim;
117- } else {
118- hcclContext_ = (__gm__ Mc2Kernel::HcclOpParam*)GetHcclContext<HCCL_GROUP_ID_0>();
119- rankIdHccl_ = Mc2Kernel::GetRankId(hcclContext_);
120- rankDimHccl_ = Mc2Kernel::GetRankDim(hcclContext_);
121- }
122-}
123- 
124template <TemplateTypeClass>98template <TemplateTypeClass>
125__aicore__ inline void MTECommunication<TemplateType>::InitHcclContext()99__aicore__ inline void MTECommunication<TemplateType>::InitHcclContext()
126{100{
127 hcclContext_ = (__gm__ Mc2Kernel::HcclOpParam*)GetHcclContext<HCCL_GROUP_ID_0>();101 hcclContext_ = (__gm__ Mc2Kernel::HcclOpParam*)GetHcclContext<HCCL_GROUP_ID_0>();
128- rankIdHccl_ = Mc2Kernel::GetRankId(hcclContext_);
129- rankDimHccl_ = Mc2Kernel::GetRankDim(hcclContext_);
130- isMc2Context_ = false;
131}102}
132 103 
133template <TemplateTypeClass>104template <TemplateTypeClass>
@@ -152,14 +123,14 @@ __aicore__ inline void MTECommunication<TemplateType>::InitGMTensor(GM_ADDR x, G
152 // =========== Win区相关 =========== 123 // =========== Win区相关 ===========
153 // Win区OOM检测适配，告知OOM框架Win区地址和大小124 // Win区OOM检测适配，告知OOM框架Win区地址和大小
154 #if defined(ASCENDC_OOM) && ASCENDC_OOM == 1125 #if defined(ASCENDC_OOM) && ASCENDC_OOM == 1
155- for (uint64_t curRank = 0; curRank < rankDimHccl_; ++curRank) {126+ for(uint64_t curRank = 0; curRank < hcclContext_->rankDim; ++curRank) {
156 OOMCheckAddrRange(GetWinAddrGm(curRank), winSpaceGmSize);127 OOMCheckAddrRange(GetWinAddrGm(curRank), winSpaceGmSize);
157 }128 }
158 #endif129 #endif
159 130 
160 // 处理 0/1 分区 标志位131 // 处理 0/1 分区 标志位
161 uint64_t currCoreFlagOffset = 2UL * SINGLE_STATE_REGION_SIZE + aivId_ * WIN_ADDR_ALIGN; // 计算当前核的标志位在Win区的偏移132 uint64_t currCoreFlagOffset = 2UL * SINGLE_STATE_REGION_SIZE + aivId_ * WIN_ADDR_ALIGN; // 计算当前核的标志位在Win区的偏移
162- selfWinFlagGMTensor_.SetGlobalBuffer((__gm__ uint32_t*)GetWinAddrGm(rankIdHccl_, currCoreFlagOffset));133+ selfWinFlagGMTensor_.SetGlobalBuffer((__gm__ uint32_t*)GetWinAddrGm(hcclContext_->rankId, currCoreFlagOffset));
163 LocalTensor<uint32_t> winFlagLocalTensor = winFlagsBuf_.Get<uint32_t>();134 LocalTensor<uint32_t> winFlagLocalTensor = winFlagsBuf_.Get<uint32_t>();
164 DataCopy(winFlagLocalTensor, selfWinFlagGMTensor_, UB_ALIGN_BYTES / sizeof(uint32_t)); // GM -> UB135 DataCopy(winFlagLocalTensor, selfWinFlagGMTensor_, UB_ALIGN_BYTES / sizeof(uint32_t)); // GM -> UB
165 SyncFunc<AscendC::HardEvent::MTE2_S>();136 SyncFunc<AscendC::HardEvent::MTE2_S>();
@@ -170,7 +141,7 @@ __aicore__ inline void MTECommunication<TemplateType>::InitGMTensor(GM_ADDR x, G
170 141 
171 // 获取本卡地址写数据142 // 获取本卡地址写数据
172 // 通过rankId和0/1分区标志位获取本地winIn区地址对应卡的数据区域143 // 通过rankId和0/1分区标志位获取本地winIn区地址对应卡的数据区域
173- GM_ADDR localDataSpaceGm = GetWinDataAddrGm(rankIdHccl_, winBufferFlags_);144+ GM_ADDR localDataSpaceGm = GetWinDataAddrGm(hcclContext_->rankId, winBufferFlags_);
174 localWinXGMTensor_.SetGlobalBuffer((__gm__ XType*)localDataSpaceGm);145 localWinXGMTensor_.SetGlobalBuffer((__gm__ XType*)localDataSpaceGm);
175 localWinScaleGMTensor_.SetGlobalBuffer((__gm__ ScalesType*)(localDataSpaceGm + xSize)); // sclae数据跟在x后146 localWinScaleGMTensor_.SetGlobalBuffer((__gm__ ScalesType*)(localDataSpaceGm + xSize)); // sclae数据跟在x后
176}147}
@@ -188,12 +159,11 @@ __aicore__ inline void MTECommunication<TemplateType>::InitBuffer(TPipe *tPipe)
188 tPipe->InitBuffer(xOutQueue_, BUFFER_NUM, xNumPerBlock_ * sizeof(OutputType)); // 用于输出的OutPutTensor159 tPipe->InitBuffer(xOutQueue_, BUFFER_NUM, xNumPerBlock_ * sizeof(OutputType)); // 用于输出的OutPutTensor
189 tPipe->InitBuffer(winFlagsBuf_, UB_ALIGN_BYTES); // 用于读取0/1分区的标志位160 tPipe->InitBuffer(winFlagsBuf_, UB_ALIGN_BYTES); // 用于读取0/1分区的标志位
190 tPipe->InitBuffer(writeStateBuf_, UB_ALIGN_BYTES); // 状态位每一个按32B对齐161 tPipe->InitBuffer(writeStateBuf_, UB_ALIGN_BYTES); // 状态位每一个按32B对齐
191- tPipe->InitBuffer(readStateBuf_, rankDimHccl_ * UB_ALIGN_BYTES); // 每次读 rankDim 个状态位162+ tPipe->InitBuffer(readStateBuf_, hcclContext_->rankDim * UB_ALIGN_BYTES); // 每次读 rankDim 个状态位
192- tPipe->InitBuffer(stateResetBuf_, rankDimHccl_ * UB_ALIGN_BYTES); // 用于清理状态区163+ tPipe->InitBuffer(stateResetBuf_, hcclContext_->rankDim * UB_ALIGN_BYTES); // 用于清理状态区
193 164 
194 stateResetTensor_ = stateResetBuf_.Get<float>();165 stateResetTensor_ = stateResetBuf_.Get<float>();
195- Duplicate<float>(stateResetTensor_, (float)0.0,166+ Duplicate<float>(stateResetTensor_, (float)0.0, static_cast<uint32_t>(hcclContext_->rankDim * FLOAT_UB_ALIGN_NUM)); // 用于状态区清零
196- static_cast<uint32_t>(rankDimHccl_ * FLOAT_UB_ALIGN_NUM)); // 用于状态区清零
197}167}
198 168 
199/**169/**
@@ -299,7 +269,7 @@ __aicore__ inline void MTECommunication<TemplateType>::CopyDataToWin(uint64_t xS
299 }269 }
300 if constexpr (isReduceScatter) {270 if constexpr (isReduceScatter) {
301 // ReduceScatter过程，数据按卡均分，需要对卡进行遍历271 // ReduceScatter过程，数据按卡均分，需要对卡进行遍历
302- for (uint64_t curRank = 0; curRank < rankDimHccl_; ++curRank) {272+ for(uint64_t curRank = 0; curRank < hcclContext_->rankDim; ++curRank) {
303 // all2all过程，加上卡偏移273 // all2all过程，加上卡偏移
304 uint64_t curRankXOffset = curXOffset + curRank * xSliceSizeNums;274 uint64_t curRankXOffset = curXOffset + curRank * xSliceSizeNums;
305 uint64_t curRankScaleOffset = curScaleOffset + curRank * scaleSliceNums;275 uint64_t curRankScaleOffset = curScaleOffset + curRank * scaleSliceNums;
@@ -326,9 +296,9 @@ __aicore__ inline void MTECommunication<TemplateType>::CopyDataToWin(uint64_t xS
326template <TemplateTypeClass>296template <TemplateTypeClass>
327__aicore__ inline void MTECommunication<TemplateType>::WriteStatusToWin()297__aicore__ inline void MTECommunication<TemplateType>::WriteStatusToWin()
328{298{
329- uint32_t coreOffset = aivId_ * rankDimHccl_; // Win区大小为 aivNum * rankDim, 此处计算核偏移299+ uint32_t coreOffset = aivId_ * hcclContext_->rankDim; // Win区大小为 aivNum * rankDim, 此处计算核偏移
330 // 遍历每一张卡，给每一张卡都要写入状态300 // 遍历每一张卡，给每一张卡都要写入状态
331- for (uint32_t curRank = 0; curRank < rankDimHccl_; ++curRank) {301+ for (uint32_t curRank = 0; curRank < hcclContext_->rankDim; ++curRank) {
332 // 写入状态到对端，每个核写一个状态，表示自己的数据块已经写完302 // 写入状态到对端，每个核写一个状态，表示自己的数据块已经写完
333 LocalTensor<float> statusTensor = writeStateBuf_.Get<float>();303 LocalTensor<float> statusTensor = writeStateBuf_.Get<float>();
334 DataCopy<float>(statusTensor, stateResetTensor_, FLOAT_UB_ALIGN_NUM); // 先重置statusTensor数据，后面累加需要Tensor内全部数据，防止脏数据304 DataCopy<float>(statusTensor, stateResetTensor_, FLOAT_UB_ALIGN_NUM); // 先重置statusTensor数据，后面累加需要Tensor内全部数据，防止脏数据
@@ -338,7 +308,7 @@ __aicore__ inline void MTECommunication<TemplateType>::WriteStatusToWin()
338 GlobalTensor<float> stateGMTensor;308 GlobalTensor<float> stateGMTensor;
339 stateGMTensor.SetGlobalBuffer((__gm__ float*)remoteWinStateGM);309 stateGMTensor.SetGlobalBuffer((__gm__ float*)remoteWinStateGM);
340 // 不同卡上的核的状态写到相邻位置，读时可以一次读rankDim个状态, 状态区大小设计为 aivNum * ranDim310 // 不同卡上的核的状态写到相邻位置，读时可以一次读rankDim个状态, 状态区大小设计为 aivNum * ranDim
341- uint64_t curOffset = (coreOffset + rankIdHccl_) * FLOAT_UB_ALIGN_NUM; // 当前核偏移 + 卡偏移， 按32B对齐311+ uint64_t curOffset = (coreOffset + hcclContext_->rankId) * FLOAT_UB_ALIGN_NUM; // 当前核偏移 + 卡偏移， 按32B对齐
342 SyncFunc<AscendC::HardEvent::S_MTE3>();312 SyncFunc<AscendC::HardEvent::S_MTE3>();
343 DataCopy(stateGMTensor[curOffset], statusTensor, FLOAT_UB_ALIGN_NUM); // 按32B对齐拷贝313 DataCopy(stateGMTensor[curOffset], statusTensor, FLOAT_UB_ALIGN_NUM); // 按32B对齐拷贝
344 SyncFunc<AscendC::HardEvent::MTE3_S>();314 SyncFunc<AscendC::HardEvent::MTE3_S>();
@@ -355,16 +325,16 @@ __aicore__ inline void MTECommunication<TemplateType>::WriteStatusToWin()
355template <TemplateTypeClass>325template <TemplateTypeClass>
356__aicore__ inline void MTECommunication<TemplateType>::ReadStatus()326__aicore__ inline void MTECommunication<TemplateType>::ReadStatus()
357{327{
358- GM_ADDR stateGM = GetWinStatusAddrGm(rankIdHccl_, winBufferFlags_); // 获取本卡的状态区用于读取328+ GM_ADDR stateGM = GetWinStatusAddrGm(hcclContext_->rankId, winBufferFlags_); // 获取本卡的状态区用于读取
359 GlobalTensor<float> selfStatusWinTensor;329 GlobalTensor<float> selfStatusWinTensor;
360- uint32_t offset = aivId_ * rankDimHccl_ * FLOAT_UB_ALIGN_NUM; // 获取当前核所需读取状态位的头地址，状态按32B对齐330+ uint32_t offset = aivId_ * hcclContext_->rankDim * FLOAT_UB_ALIGN_NUM; // 获取当前核所需读取状态位的头地址，状态按32B对齐
361 selfStatusWinTensor.SetGlobalBuffer((__gm__ float*)(stateGM));331 selfStatusWinTensor.SetGlobalBuffer((__gm__ float*)(stateGM));
362 LocalTensor<float> statusTensor = readStateBuf_.Get<float>();332 LocalTensor<float> statusTensor = readStateBuf_.Get<float>();
363 float flag = 0; // 用于计算状态和333 float flag = 0; // 用于计算状态和
364- uint32_t statusCnt = rankDimHccl_ * FLOAT_UB_ALIGN_NUM; // 一次读rankDim个，按32B对齐334+ uint32_t statusCnt = hcclContext_->rankDim * FLOAT_UB_ALIGN_NUM; // 一次读rankDim个，按32B对齐
365 SumParams sumParams{1, statusCnt, statusCnt};335 SumParams sumParams{1, statusCnt, statusCnt};
366- float minTarget = rankDimHccl_ - (float)0.5;336+ float minTarget = hcclContext_->rankDim - (float)0.5;
367- float maxTarget = rankDimHccl_ + (float)0.5;337+ float maxTarget = hcclContext_->rankDim + (float)0.5;
368 // 读取statusCnt个数据求和338 // 读取statusCnt个数据求和
369 while ((flag < minTarget) || (flag > maxTarget)) {339 while ((flag < minTarget) || (flag > maxTarget)) {
370 SyncFunc<AscendC::HardEvent::S_MTE2>();340 SyncFunc<AscendC::HardEvent::S_MTE2>();
@@ -405,9 +375,6 @@ __aicore__ inline void MTECommunication<TemplateType>::CopyResultToOutput(uint64
405template <TemplateTypeClass>375template <TemplateTypeClass>
406__aicore__ inline GM_ADDR MTECommunication<TemplateType>::GetWinAddrGm(uint32_t rankId, uint64_t offset)376__aicore__ inline GM_ADDR MTECommunication<TemplateType>::GetWinAddrGm(uint32_t rankId, uint64_t offset)
407{377{
408- if (isMc2Context_) {
409- return (GM_ADDR)(mc2Context_->windowsIn[rankId] + offset);
410- }
411 return (GM_ADDR)(hcclContext_->windowsIn[rankId] + offset);378 return (GM_ADDR)(hcclContext_->windowsIn[rankId] + offset);
412}379}
413 380 
@@ -420,9 +387,6 @@ __aicore__ inline GM_ADDR MTECommunication<TemplateType>::GetWinDataAddrGm(uint3
420 return GetWinAddrGm(rankId, STATE_WIN_SIZE);387 return GetWinAddrGm(rankId, STATE_WIN_SIZE);
421 }388 }
422 else {389 else {
423- if (isMc2Context_) {
424- return (GM_ADDR)(mc2Context_->windowsOut[rankId]);
425- }
426 // 若使用 1 分区，即WinOut390 // 若使用 1 分区，即WinOut
427 return (GM_ADDR)(hcclContext_->windowsOut[rankId]);391 return (GM_ADDR)(hcclContext_->windowsOut[rankId]);
428 }392 }
@@ -441,7 +405,5 @@ __aicore__ inline GM_ADDR MTECommunication<TemplateType>::GetWinStatusAddrGm(uin
441 return GetWinAddrGm(rankId, SINGLE_STATE_REGION_SIZE);405 return GetWinAddrGm(rankId, SINGLE_STATE_REGION_SIZE);
442 }406 }
443}407}
444- 
445} // QuantMTECommImpl408} // QuantMTECommImpl
446- 
447#endif // MTE_COMMON_H409#endif // MTE_COMMON_H
Mmc2/quant_reduce_scatter/op_kernel/quant_reduce_scatter.cpp+2-2
@@ -29,7 +29,7 @@ using namespace QuantReduceScatterImpl;
29 29 
30template<uint32_t quantReduceScatterCommMode>30template<uint32_t quantReduceScatterCommMode>
31__global__ __aicore__ void quant_reduce_scatter(GM_ADDR x, GM_ADDR scales, GM_ADDR output, GM_ADDR workspaceGM,31__global__ __aicore__ void quant_reduce_scatter(GM_ADDR x, GM_ADDR scales, GM_ADDR output, GM_ADDR workspaceGM,
32- GM_ADDR tilingGM)32+ GM_ADDR tilingGM)
33{33{
34 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);34 KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
35 REGISTER_TILING_DEFAULT(QuantReduceScatterTilingData);35 REGISTER_TILING_DEFAULT(QuantReduceScatterTilingData);
@@ -37,7 +37,7 @@ __global__ __aicore__ void quant_reduce_scatter(GM_ADDR x, GM_ADDR scales, GM_AD
37 TPipe pipe;37 TPipe pipe;
38 if constexpr (quantReduceScatterCommMode == MTE_COMM) {38 if constexpr (quantReduceScatterCommMode == MTE_COMM) {
39 QuantReduceScatterMte<DTYPE_X, DTYPE_SCALES, DTYPE_OUT_PUT> op;39 QuantReduceScatterMte<DTYPE_X, DTYPE_SCALES, DTYPE_OUT_PUT> op;
40- op.Init(nullptr, x, scales, output, &pipe, &tilingData);40+ op.Init(x, scales, output, &pipe, &tilingData);
41 op.Process();41 op.Process();
42 }42 }
43}43}
Mmc2/quant_reduce_scatter/op_kernel/quant_reduce_scatter_mte.h+13-13
@@ -35,7 +35,6 @@ namespace QuantReduceScatterImpl {
35using namespace QuantMTECommImpl;35using namespace QuantMTECommImpl;
36using namespace VectorComputeImpl;36using namespace VectorComputeImpl;
37using namespace AscendC;37using namespace AscendC;
38-using namespace Mc2Aclnn;
39 38 
40constexpr static uint64_t MX_SCALES_LAST_DIM = 2U; // MX量化scales最后一维的大小39constexpr static uint64_t MX_SCALES_LAST_DIM = 2U; // MX量化scales最后一维的大小
41 40 
@@ -43,7 +42,7 @@ template<TemplateTypeClass>
43class QuantReduceScatterMte {42class QuantReduceScatterMte {
44public:43public:
45 __aicore__ inline QuantReduceScatterMte() {};44 __aicore__ inline QuantReduceScatterMte() {};
46- __aicore__ inline void Init(GM_ADDR mc2Context, GM_ADDR x, GM_ADDR scales, GM_ADDR output,45+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR scales, GM_ADDR output,
47 TPipe *pipe, const QuantReduceScatterTilingData *tilingData);46 TPipe *pipe, const QuantReduceScatterTilingData *tilingData);
48 __aicore__ inline void Process();47 __aicore__ inline void Process();
49private:48private:
@@ -87,10 +86,10 @@ private:
87};86};
88 87 
89template <TemplateTypeClass>88template <TemplateTypeClass>
90-__aicore__ inline void QuantReduceScatterMte<TemplateType>::Init(GM_ADDR mc2Context, GM_ADDR x, GM_ADDR scales,89+__aicore__ inline void QuantReduceScatterMte<TemplateType>::Init(GM_ADDR x, GM_ADDR scales,
91 GM_ADDR output, TPipe *tPipe, const QuantReduceScatterTilingData *tilingData)90 GM_ADDR output, TPipe *tPipe, const QuantReduceScatterTilingData *tilingData)
92{91{
93- mteComm_.InitMc2Context(mc2Context, tilingData);92+ mteComm_.InitHcclContext();
94 ParseTilingInfo(tilingData);93 ParseTilingInfo(tilingData);
95 tPipe->Reset();94 tPipe->Reset();
96 ComputeXPerBlock(tilingData, tPipe);95 ComputeXPerBlock(tilingData, tPipe);
@@ -110,9 +109,9 @@ __aicore__ inline void QuantReduceScatterMte<TemplateType>::ParseTilingInfo(
110 if constexpr(AscendC::IsSameType<ScalesType, fp8_e8m0_t>::value) {109 if constexpr(AscendC::IsSameType<ScalesType, fp8_e8m0_t>::value) {
111 scaleSize_ *= MX_SCALES_LAST_DIM;110 scaleSize_ *= MX_SCALES_LAST_DIM;
112 }111 }
113- uint64_t xSliceSize = xSize_ / (mteComm_.rankDimHccl_);112+ uint64_t xSliceSize = xSize_ / (mteComm_.hcclContext_->rankDim);
114 xSliceSizeNums_ = xSliceSize / sizeof(XType);113 xSliceSizeNums_ = xSliceSize / sizeof(XType);
115- scaleSliceNums_ = scaleSize_ / (mteComm_.rankDimHccl_ * sizeof(ScalesType));114+ scaleSliceNums_ = scaleSize_ / (mteComm_.hcclContext_->rankDim * sizeof(ScalesType));
116}115}
117 116 
118template <TemplateTypeClass>117template <TemplateTypeClass>
@@ -125,8 +124,8 @@ __aicore__ inline void QuantReduceScatterMte<TemplateType>::ComputeXPerBlock(
125 uint64_t mteCommFixedSpace = BUFFER_NUM * X_BLOCK_BYTES + // scaleQueue_124 uint64_t mteCommFixedSpace = BUFFER_NUM * X_BLOCK_BYTES + // scaleQueue_
126 UB_ALIGN_BYTES + // winFlagsBuf_125 UB_ALIGN_BYTES + // winFlagsBuf_
127 UB_ALIGN_BYTES + // writeStateBuf_126 UB_ALIGN_BYTES + // writeStateBuf_
128- mteComm_.rankDimHccl_ * UB_ALIGN_BYTES + // readStateBuf_127+ mteComm_.hcclContext_->rankDim * UB_ALIGN_BYTES + // readStateBuf_
129- mteComm_.rankDimHccl_ * UB_ALIGN_BYTES; // stateResetBuf_128+ mteComm_.hcclContext_->rankDim * UB_ALIGN_BYTES; // stateResetBuf_
130 129 
131 // 动态开销：每增加 1 个 x 需要的 UB 字节（整数部分，分数部分见下方比例校正）130 // 动态开销：每增加 1 个 x 需要的 UB 字节（整数部分，分数部分见下方比例校正）
132 uint64_t baseDynamic = BUFFER_NUM * sizeof(OutputType) + // xOutQueue_131 uint64_t baseDynamic = BUFFER_NUM * sizeof(OutputType) + // xOutQueue_
@@ -240,6 +239,7 @@ __aicore__ inline void QuantReduceScatterMte<TemplateType>::ReadDataBlockReduceS
240 scaleInQue_.FreeTensor(scaleTmpTensor);239 scaleInQue_.FreeTensor(scaleTmpTensor);
241}240}
242 241 
242+ 
243template <TemplateTypeClass>243template <TemplateTypeClass>
244__aicore__ inline void QuantReduceScatterMte<TemplateType>::ClearSumTensor()244__aicore__ inline void QuantReduceScatterMte<TemplateType>::ClearSumTensor()
245{245{
@@ -275,9 +275,9 @@ __aicore__ inline void QuantReduceScatterMte<TemplateType>::ExecuteReduceScatter
275 275 
276 // 遍历每张卡，读取其Win区的数据，采取错卡序读取，从自己卡上读起276 // 遍历每张卡，读取其Win区的数据，采取错卡序读取，从自己卡上读起
277 /* rank0: [0,1,2]; rank1: [1,2,0]; rank2: [2,0,1] */277 /* rank0: [0,1,2]; rank1: [1,2,0]; rank2: [2,0,1] */
278- uint32_t startRankId = mteComm_.rankIdHccl_;278+ uint32_t startRankId = mteComm_.hcclContext_->rankId;
279- for (uint32_t i = 0; i < mteComm_.rankDimHccl_; ++i) {279+ for (uint32_t i = 0; i < mteComm_.hcclContext_->rankDim; ++i) {
280- uint32_t remoteRankId = (startRankId + i) % mteComm_.rankDimHccl_;280+ uint32_t remoteRankId = (startRankId + i) % mteComm_.hcclContext_->rankDim;
281 281 
282 // 获取对端Win区中数据区相关的地址282 // 获取对端Win区中数据区相关的地址
283 GM_ADDR remoteDataSpaceGm = mteComm_.GetWinDataAddrGm(remoteRankId, mteComm_.winBufferFlags_);283 GM_ADDR remoteDataSpaceGm = mteComm_.GetWinDataAddrGm(remoteRankId, mteComm_.winBufferFlags_);
@@ -287,8 +287,8 @@ __aicore__ inline void QuantReduceScatterMte<TemplateType>::ExecuteReduceScatter
287 287 
288 // 读取对端对应地址的 x 和 scale数据，进行反量化和求和288 // 读取对端对应地址的 x 和 scale数据，进行反量化和求和
289 // ReduceScatter过程，all2all仅需与rankId相关的数据，加上本卡偏移289 // ReduceScatter过程，all2all仅需与rankId相关的数据，加上本卡偏移
290- uint64_t curRankXOffset = curXOffset + mteComm_.rankIdHccl_ * xSliceSizeNums_;290+ uint64_t curRankXOffset = curXOffset + mteComm_.hcclContext_->rankId * xSliceSizeNums_;
291- uint64_t curRankScaleOffset = curScaleOffset + mteComm_.rankIdHccl_ * scaleSliceNums_;291+ uint64_t curRankScaleOffset = curScaleOffset + mteComm_.hcclContext_->rankId * scaleSliceNums_;
292 ReadDataBlockReduceSum(curRankXOffset, curRankScaleOffset, curXNum, curScaleNum);292 ReadDataBlockReduceSum(curRankXOffset, curRankScaleOffset, curXNum, curScaleNum);
293 }293 }
294 294 
Mmc2/quant_reduce_scatter/op_kernel/quant_reduce_scatter_tiling_data.h+0-1
@@ -27,7 +27,6 @@ struct QuantReduceScatterTilingInfo {
27 uint64_t totalWinSize; // Win区总大小，即HCCL_BUFFER_SIZE27 uint64_t totalWinSize; // Win区总大小，即HCCL_BUFFER_SIZE
28 uint32_t xPerBlock; // host 侧基于 TARGET_ITER 公式推荐的每块元素数28 uint32_t xPerBlock; // host 侧基于 TARGET_ITER 公式推荐的每块元素数
29 uint32_t alignBlock; // xPerBlock 对齐粒度（元素数，host/kernel共享）29 uint32_t alignBlock; // xPerBlock 对齐粒度（元素数，host/kernel共享）
30- bool isMc2Context;
31};30};
32 31 
33struct QuantReduceScatterTilingData {32struct QuantReduceScatterTilingData {
Mmc2/quant_reduce_scatter/tests/ut/op_host/arch35/test_quant_reduce_scatter_tiling.cpp+14-14
@@ -65,37 +65,37 @@ static QuantReduceScatterTestParam g_testCases[] = {
65 {1024, 80, 2}, ge::DT_FLOAT8_E8M0, ge::FORMAT_ND,65 {1024, 80, 2}, ge::DT_FLOAT8_E8M0, ge::FORMAT_ND,
66 {128, 5120}, ge::DT_FLOAT16, ge::FORMAT_ND,66 {128, 5120}, ge::DT_FLOAT16, ge::FORMAT_ND,
67 "group", "sum", ge::DT_FLOAT16, 8, "3510",67 "group", "sum", ge::DT_FLOAT16, 8, "3510",
68- ge::GRAPH_SUCCESS, 0UL, "1024 5120 80 64 314572800 4398046514176 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},68+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 80 64 314572800 4398046514176 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
69 {"quant_reduce_scatter_critical_case_2",69 {"quant_reduce_scatter_critical_case_2",
70 {2048, 5120}, ge::DT_HIFLOAT8, ge::FORMAT_ND,70 {2048, 5120}, ge::DT_HIFLOAT8, ge::FORMAT_ND,
71 {2048, 40}, ge::DT_FLOAT, ge::FORMAT_ND,71 {2048, 40}, ge::DT_FLOAT, ge::FORMAT_ND,
72 {256, 5120}, ge::DT_FLOAT, ge::FORMAT_ND,72 {256, 5120}, ge::DT_FLOAT, ge::FORMAT_ND,
73 "group", "sum", ge::DT_FLOAT, 8, "3510",73 "group", "sum", ge::DT_FLOAT, 8, "3510",
74- ge::GRAPH_SUCCESS, 0UL, "2048 5120 40 64 314572800 4398046517248 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},74+ ge::GRAPH_SUCCESS, 0UL, "2048 5120 40 64 314572800 4398046517248 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
75 {"quant_reduce_scatter_critical_case_3",75 {"quant_reduce_scatter_critical_case_3",
76 {1024, 7168}, ge::DT_FLOAT8_E5M2, ge::FORMAT_ND,76 {1024, 7168}, ge::DT_FLOAT8_E5M2, ge::FORMAT_ND,
77 {1024, 56}, ge::DT_FLOAT, ge::FORMAT_ND,77 {1024, 56}, ge::DT_FLOAT, ge::FORMAT_ND,
78 {128, 7168}, ge::DT_FLOAT16, ge::FORMAT_ND,78 {128, 7168}, ge::DT_FLOAT16, ge::FORMAT_ND,
79 "group", "sum", ge::DT_FLOAT16, 8, "3510",79 "group", "sum", ge::DT_FLOAT16, 8, "3510",
80- ge::GRAPH_SUCCESS, 0UL, "1024 7168 56 64 314572800 4398046515200 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},80+ ge::GRAPH_SUCCESS, 0UL, "1024 7168 56 64 314572800 4398046515200 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
81 {"quant_reduce_scatter_critical_case_1_x3d",81 {"quant_reduce_scatter_critical_case_1_x3d",
82 {8, 128, 4096}, ge::DT_FLOAT8_E4M3FN, ge::FORMAT_ND,82 {8, 128, 4096}, ge::DT_FLOAT8_E4M3FN, ge::FORMAT_ND,
83 {8, 128, 64, 2}, ge::DT_FLOAT8_E8M0, ge::FORMAT_ND,83 {8, 128, 64, 2}, ge::DT_FLOAT8_E8M0, ge::FORMAT_ND,
84 {128, 4096}, ge::DT_FLOAT16, ge::FORMAT_ND,84 {128, 4096}, ge::DT_FLOAT16, ge::FORMAT_ND,
85 "group", "sum", ge::DT_FLOAT16, 8, "3510",85 "group", "sum", ge::DT_FLOAT16, 8, "3510",
86- ge::GRAPH_SUCCESS, 0UL, "1024 4096 64 64 314572800 4398046513152 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},86+ ge::GRAPH_SUCCESS, 0UL, "1024 4096 64 64 314572800 4398046513152 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
87 {"quant_reduce_scatter_critical_case_2_x3d",87 {"quant_reduce_scatter_critical_case_2_x3d",
88 {16, 128, 4096}, ge::DT_HIFLOAT8, ge::FORMAT_ND,88 {16, 128, 4096}, ge::DT_HIFLOAT8, ge::FORMAT_ND,
89 {16, 128, 32}, ge::DT_FLOAT, ge::FORMAT_ND,89 {16, 128, 32}, ge::DT_FLOAT, ge::FORMAT_ND,
90 {256, 4096}, ge::DT_FLOAT, ge::FORMAT_ND,90 {256, 4096}, ge::DT_FLOAT, ge::FORMAT_ND,
91 "group", "sum", ge::DT_FLOAT, 8, "3510",91 "group", "sum", ge::DT_FLOAT, 8, "3510",
92- ge::GRAPH_SUCCESS, 0UL, "2048 4096 32 64 314572800 4398046516224 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},92+ ge::GRAPH_SUCCESS, 0UL, "2048 4096 32 64 314572800 4398046516224 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
93 {"quant_reduce_scatter_critical_case_3_x3d",93 {"quant_reduce_scatter_critical_case_3_x3d",
94 {8, 128, 8192}, ge::DT_FLOAT8_E5M2, ge::FORMAT_ND,94 {8, 128, 8192}, ge::DT_FLOAT8_E5M2, ge::FORMAT_ND,
95 {8, 128, 64}, ge::DT_FLOAT, ge::FORMAT_ND,95 {8, 128, 64}, ge::DT_FLOAT, ge::FORMAT_ND,
96 {128, 8192}, ge::DT_FLOAT16, ge::FORMAT_ND,96 {128, 8192}, ge::DT_FLOAT16, ge::FORMAT_ND,
97 "group", "sum", ge::DT_FLOAT16, 8, "3510",97 "group", "sum", ge::DT_FLOAT16, 8, "3510",
98- ge::GRAPH_SUCCESS, 0UL, "1024 8192 64 64 314572800 4398046516224 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},98+ ge::GRAPH_SUCCESS, 0UL, "1024 8192 64 64 314572800 4398046516224 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
99 {"quant_reduce_scatter_abuse_case_1_x3d",99 {"quant_reduce_scatter_abuse_case_1_x3d",
100 {8, 128, 8192}, ge::DT_FLOAT8_E5M2, ge::FORMAT_ND,100 {8, 128, 8192}, ge::DT_FLOAT8_E5M2, ge::FORMAT_ND,
101 {8, 128, 64}, ge::DT_FLOAT, ge::FORMAT_ND,101 {8, 128, 64}, ge::DT_FLOAT, ge::FORMAT_ND,
@@ -264,49 +264,49 @@ static QuantReduceScatterTestParam g_testCases[] = {
264 {1024, 80, 2}, ge::DT_FLOAT8_E8M0, ge::FORMAT_ND,264 {1024, 80, 2}, ge::DT_FLOAT8_E8M0, ge::FORMAT_ND,
265 {512, 5120}, ge::DT_FLOAT16, ge::FORMAT_ND,265 {512, 5120}, ge::DT_FLOAT16, ge::FORMAT_ND,
266 "group", "sum", ge::DT_FLOAT16, 2, "3510",266 "group", "sum", ge::DT_FLOAT16, 2, "3510",
267- ge::GRAPH_SUCCESS, 0UL, "1024 5120 80 64 314572800 4398046524416 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},267+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 80 64 314572800 4398046524416 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
268 {"quant_reduce_scatter_critical_case_rank4_tg",268 {"quant_reduce_scatter_critical_case_rank4_tg",
269 {1024, 5120}, ge::DT_INT8, ge::FORMAT_ND,269 {1024, 5120}, ge::DT_INT8, ge::FORMAT_ND,
270 {1024, 40}, ge::DT_FLOAT, ge::FORMAT_ND,270 {1024, 40}, ge::DT_FLOAT, ge::FORMAT_ND,
271 {256, 5120}, ge::DT_FLOAT16, ge::FORMAT_ND,271 {256, 5120}, ge::DT_FLOAT16, ge::FORMAT_ND,
272 "group", "sum", ge::DT_FLOAT16, 4, "3510",272 "group", "sum", ge::DT_FLOAT16, 4, "3510",
273- ge::GRAPH_SUCCESS, 0UL, "1024 5120 40 64 314572800 4398046517248 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},273+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 40 64 314572800 4398046517248 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
274 {"quant_reduce_scatter_critical_case_H1024_tg",274 {"quant_reduce_scatter_critical_case_H1024_tg",
275 {1024, 1024}, ge::DT_INT8, ge::FORMAT_ND,275 {1024, 1024}, ge::DT_INT8, ge::FORMAT_ND,
276 {1024, 8}, ge::DT_FLOAT, ge::FORMAT_ND,276 {1024, 8}, ge::DT_FLOAT, ge::FORMAT_ND,
277 {128, 1024}, ge::DT_FLOAT16, ge::FORMAT_ND,277 {128, 1024}, ge::DT_FLOAT16, ge::FORMAT_ND,
278 "group", "sum", ge::DT_FLOAT16, 8, "3510",278 "group", "sum", ge::DT_FLOAT16, 8, "3510",
279- ge::GRAPH_SUCCESS, 0UL, "1024 1024 8 64 314572800 4398046513152 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},279+ ge::GRAPH_SUCCESS, 0UL, "1024 1024 8 64 314572800 4398046513152 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
280 {"quant_reduce_scatter_critical_case_H8192_mx",280 {"quant_reduce_scatter_critical_case_H8192_mx",
281 {1024, 8192}, ge::DT_FLOAT8_E4M3FN, ge::FORMAT_ND,281 {1024, 8192}, ge::DT_FLOAT8_E4M3FN, ge::FORMAT_ND,
282 {1024, 128, 2}, ge::DT_FLOAT8_E8M0, ge::FORMAT_ND,282 {1024, 128, 2}, ge::DT_FLOAT8_E8M0, ge::FORMAT_ND,
283 {128, 8192}, ge::DT_FLOAT16, ge::FORMAT_ND,283 {128, 8192}, ge::DT_FLOAT16, ge::FORMAT_ND,
284 "group", "sum", ge::DT_FLOAT16, 8, "3510",284 "group", "sum", ge::DT_FLOAT16, 8, "3510",
285- ge::GRAPH_SUCCESS, 0UL, "1024 8192 128 64 314572800 4398046516224 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},285+ ge::GRAPH_SUCCESS, 0UL, "1024 8192 128 64 314572800 4398046516224 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
286 {"quant_reduce_scatter_critical_case_3d_mx_e5m2",286 {"quant_reduce_scatter_critical_case_3d_mx_e5m2",
287 {4, 256, 5120}, ge::DT_FLOAT8_E5M2, ge::FORMAT_ND,287 {4, 256, 5120}, ge::DT_FLOAT8_E5M2, ge::FORMAT_ND,
288 {4, 256, 80, 2}, ge::DT_FLOAT8_E8M0, ge::FORMAT_ND,288 {4, 256, 80, 2}, ge::DT_FLOAT8_E8M0, ge::FORMAT_ND,
289 {128, 5120}, ge::DT_BF16, ge::FORMAT_ND,289 {128, 5120}, ge::DT_BF16, ge::FORMAT_ND,
290 "group", "sum", ge::DT_BF16, 8, "3510",290 "group", "sum", ge::DT_BF16, 8, "3510",
291- ge::GRAPH_SUCCESS, 0UL, "1024 5120 80 64 314572800 4398046514176 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},291+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 80 64 314572800 4398046514176 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
292 {"quant_reduce_scatter_critical_case_3d_tg_int8_float_out",292 {"quant_reduce_scatter_critical_case_3d_tg_int8_float_out",
293 {8, 128, 5120}, ge::DT_INT8, ge::FORMAT_ND,293 {8, 128, 5120}, ge::DT_INT8, ge::FORMAT_ND,
294 {8, 128, 40}, ge::DT_FLOAT, ge::FORMAT_ND,294 {8, 128, 40}, ge::DT_FLOAT, ge::FORMAT_ND,
295 {128, 5120}, ge::DT_FLOAT, ge::FORMAT_ND,295 {128, 5120}, ge::DT_FLOAT, ge::FORMAT_ND,
296 "group", "sum", ge::DT_FLOAT, 8, "3510",296 "group", "sum", ge::DT_FLOAT, 8, "3510",
297- ge::GRAPH_SUCCESS, 0UL, "1024 5120 40 64 314572800 4398046514176 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},297+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 40 64 314572800 4398046514176 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
298 {"quant_reduce_scatter_critical_case_hifloat8_bf16",298 {"quant_reduce_scatter_critical_case_hifloat8_bf16",
299 {1024, 5120}, ge::DT_HIFLOAT8, ge::FORMAT_ND,299 {1024, 5120}, ge::DT_HIFLOAT8, ge::FORMAT_ND,
300 {1024, 40}, ge::DT_FLOAT, ge::FORMAT_ND,300 {1024, 40}, ge::DT_FLOAT, ge::FORMAT_ND,
301 {128, 5120}, ge::DT_BF16, ge::FORMAT_ND,301 {128, 5120}, ge::DT_BF16, ge::FORMAT_ND,
302 "group", "sum", ge::DT_BF16, 8, "3510",302 "group", "sum", ge::DT_BF16, 8, "3510",
303- ge::GRAPH_SUCCESS, 0UL, "1024 5120 40 64 314572800 4398046514176 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},303+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 40 64 314572800 4398046514176 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
304 {"quant_reduce_scatter_critical_case_3d_tg_rank2",304 {"quant_reduce_scatter_critical_case_3d_tg_rank2",
305 {2, 512, 5120}, ge::DT_INT8, ge::FORMAT_ND,305 {2, 512, 5120}, ge::DT_INT8, ge::FORMAT_ND,
306 {2, 512, 40}, ge::DT_FLOAT, ge::FORMAT_ND,306 {2, 512, 40}, ge::DT_FLOAT, ge::FORMAT_ND,
307 {512, 5120}, ge::DT_FLOAT16, ge::FORMAT_ND,307 {512, 5120}, ge::DT_FLOAT16, ge::FORMAT_ND,
308 "group", "sum", ge::DT_FLOAT16, 2, "3510",308 "group", "sum", ge::DT_FLOAT16, 2, "3510",
309- ge::GRAPH_SUCCESS, 0UL, "1024 5120 40 64 314572800 4398046524416 0 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},309+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 40 64 314572800 4398046524416 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},
310 // --- 新增异常路径用例 ---310 // --- 新增异常路径用例 ---
311 {"quant_reduce_scatter_abuse_case_1d_x",311 {"quant_reduce_scatter_abuse_case_1d_x",
312 {5120}, ge::DT_INT8, ge::FORMAT_ND,312 {5120}, ge::DT_INT8, ge::FORMAT_ND,
Dmc2/quant_reduce_scatter_v2/CMakeLists.txt+0-21
@@ -1,21 +0,0 @@
1-# -----------------------------------------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-# CANN Open Software License Agreement Version 2.0 (the "License").
5-# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-# See LICENSE in the root of the software repository for the full text of the License.
9-# -----------------------------------------------------------------------------------------------------------
10- 
11-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12-if(NOT ENABLE_TEST)
13- list(REMOVE_ITEM CURRENT_DIRS tests)
14-endif()
15-foreach(SUB_DIR ${CURRENT_DIRS})
16- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
17- add_subdirectory(${SUB_DIR})
18- endif()
19-endforeach()
20- 
21-set(MC2_COMPILE ${SUB_MC2_COMPILE} PARENT_SCOPE)
Dmc2/quant_reduce_scatter_v2/op_graph/CMakeLists.txt+0-14
@@ -1,14 +0,0 @@
1-# -----------------------------------------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-# CANN Open Software License Agreement Version 2.0 (the "License").
5-# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-# See LICENSE in the root of the software repository for the full text of the License.
9-# -----------------------------------------------------------------------------------------------------------
10- 
11-message(STATUS "=== Debug: start ops.transformer.quant_reduce_scatter_v2.CMakeLists.txt ")
12-if (BUILD_OPEN_PROJECT)
13- add_graph_plugin_sources()
14-endif()
Dmc2/quant_reduce_scatter_v2/op_graph/quant_reduce_scatter_v2_proto.h+0-58
@@ -1,58 +0,0 @@
1-/**
2- * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4- * CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9-**/
10- 
11-/*!
12- * \file quant_reduce_scatter_v2_proto.h
13- * \brief 图模式原型定义
14- */
15-#ifndef QUANT_REDUCE_SCATTER_V2_PROTO_H_
16-#define QUANT_REDUCE_SCATTER_V2_PROTO_H_
17- 
18-#include <graph/operator_reg.h>
19- 
20-namespace ge {
21- 
22-/**
23- * @brief Fusion op of quant and reduce scatter.
24- * @par Inputs:
25- * two inputs, including:
26- * @li context: A tensor. Support dtype: int32, dimension must be 1, Support Shape (272, ), support format: ND.
27- * @li x: A matrix tensor. The type support int8, hifloat8, float8_e4m3fn, float8_e5m2, float4_e1m2, float4_e2m1.
28- * The format supports ND.
29- * @li scale: A matrix tensor. The type support float32, float8_e8m0. The format supports ND.
30- *
31- * @par Outputs:
32- * out_put: A matrix tensor. The type support float16, bfloat16, float32. The format supports ND.
33- *
34- * @par Attributes:
35- * @li group: A required string identifying the group of ranks participating in the op.
36- * @li reduce_op: An optional string identifying the reduction operation to perform. Default: "sum".
37- * @li output_dtype: An optional int identifying the data type of output.
38- * The type support 0(float), 1(float16), 27(bfloat16). Default: 27(bfloat16).
39- * @li world_size: A required int identifying the rank size.
40- */
41-REG_OP(QuantReduceScatterV2)
42- .INPUT(context, "T0")
43- .INPUT(x, "T1")
44- .INPUT(scales, "T2")
45- .OUTPUT(out_put, "T3")
46- .DATATYPE(T0, TensorType({DT_INT32}))
47- .DATATYPE(T1, TensorType({DT_INT8, DT_HIFLOAT8, DT_FLOAT8_E5M2, DT_FLOAT8_E4M3FN, DT_FLOAT4_E1M2, DT_FLOAT4_E2M1}))
48- .DATATYPE(T2, TensorType({DT_FLOAT, DT_FLOAT8_E8M0}))
49- .DATATYPE(T3, TensorType({DT_FLOAT16, DT_BF16, DT_FLOAT}))
50- .REQUIRED_ATTR(hccl_buffer_size, Int)
51- .ATTR(reduce_op, String, "sum")
52- .ATTR(output_dtype, Int, DT_BF16)
53- .REQUIRED_ATTR(world_size, Int)
54- .OP_END_FACTORY_REG(QuantReduceScatterV2)
55- 
56-} // namespace ge
57- 
58-#endif // QUANT_REDUCE_SCATTER_V2_PROTO_H_
Dmc2/quant_reduce_scatter_v2/op_host/CMakeLists.txt+0-35
@@ -1,35 +0,0 @@
1-# -----------------------------------------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-# CANN Open Software License Agreement Version 2.0 (the "License").
5-# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-# See LICENSE in the root of the software repository for the full text of the License.
9-# -----------------------------------------------------------------------------------------------------------
10- 
11-if (BUILD_OPEN_PROJECT) # custom
12- target_sources(op_host_aclnnInner PRIVATE
13- quant_reduce_scatter_v2_def.cpp
14- )
15- add_modules_sources_with_soc(
16- OP_API_INDEPENDENT ON
17- OP_API_DIR ${CMAKE_CURRENT_SOURCE_DIR}/../op_api
18- OP_MC2_ENABLE ON
19- OPTYPE quant_reduce_scatter_v2 ACLNNTYPE aclnn_inner)
20- set(SUB_MC2_COMPILE TRUE PARENT_SCOPE)
21- set(MC2_OPT ON PARENT_SCOPE)
22- set(quant_reduce_scatter_v2_depends mc2/common mc2/3rd mc2/quant_reduce_scatter PARENT_SCOPE)
23- 
24- # --cce-auto-sync=off：指定CCE编译器是否自动执行线程间或模块间的同步操作
25- set(CONDITION_UNIT ${ASCEND_COMPUTE_UNIT})
26- if("${CONDITION_UNIT}" STREQUAL "ascend950")
27- add_ops_compile_options(
28- OP_NAME QuantReduceScatterV2
29- OPTIONS --cce-auto-sync=off
30- )
31- endif()
32-else() # 回黄host
33- add_mc2_modules_sources(OPTYPE quant_reduce_scatter_v2 ACLNNTYPE aclnn_inner)
34- set(SUB_MC2_COMPILE TRUE PARENT_SCOPE)
35-endif()
Dmc2/quant_reduce_scatter_v2/op_host/op_tiling/quant_reduce_scatter_v2_tiling.cpp+0-171
@@ -1,171 +0,0 @@
1-/**
2- * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4- * CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9-**/
10- 
11-/*!
12- * \file quant_reduce_scatter_v2_tiling.cpp
13- * \brief host侧tiling实现
14- */
15-#include <register/op_def_registry.h>
16-#include "mc2/quant_reduce_scatter/op_kernel/quant_reduce_scatter_tiling_key.h"
17-#include "mc2/quant_reduce_scatter/op_kernel/quant_reduce_scatter_tiling_data.h"
18-#include "mc2/quant_reduce_scatter/op_host/op_tiling/common/quant_reduce_scatter_util_tiling.h"
19- 
20-namespace MC2Tiling {
21- 
22-using namespace AscendC;
23-using namespace ge;
24- 
25-/**
26- * @brief 设置tilingData，给各成员变量赋值
27- */
28-static void SetTilingData(gert::TilingContext *context, QuantReduceScatterTilingData &tilingData,
29- const QuantReduceScatterConfig& config)
30-{
31- fe::PlatFormInfos *platformInfoPtr = context->GetPlatformInfo();
32- platform_ascendc::PlatformAscendC ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
33- // set tilingData
34- uint32_t aivNum = ascendcPlatform.GetCoreNumAiv();
35- context->SetBlockDim(ascendcPlatform.CalcTschBlockDim(aivNum, 0, aivNum));
36- tilingData.quantReduceScatterTilingInfo.aivNum = aivNum;
37- uint64_t xValueBS = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);
38- uint64_t xValueH = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
39- uint64_t scalesValueH = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_ONE);
40- // 3d场景，context->GetInputShape在函数CheckInputTensorDim中已经校验
41- if (context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDimNum() == THREE_DIMS) {
42- xValueBS = xValueBS * context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
43- xValueH = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_TWO);
44- scalesValueH = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_TWO);
45- }
46- tilingData.quantReduceScatterTilingInfo.bs = xValueBS;
47- tilingData.quantReduceScatterTilingInfo.hiddenSize = xValueH;
48- tilingData.quantReduceScatterTilingInfo.scaleHiddenSize = scalesValueH;
49- tilingData.quantReduceScatterTilingInfo.totalWinSize = mc2tiling::Mc2TilingUtils::GetMaxWindowSize();
50- tilingData.quantReduceScatterTilingInfo.isMc2Context = config.isMc2Context;
51-}
52- 
53-/**
54- * @brief 基于 TARGET_ITER 公式计算 host 推荐的 xPerBlock（先除 rankSize 再反推），写入 tilingData
55- */
56-static void SetXPerBlock(QuantReduceScatterTilingData &tilingData, const TilingRunInfo &runInfo)
57-{
58- constexpr uint32_t TARGET_ITER = 3U; // 与 QAR 对称
59- constexpr uint32_t MIN_BLOCK = 2048U; // QRS comm-bound：per_core 太小公式失效，MIN 兜到 2048
60- constexpr uint32_t ALIGN_BLOCK = 1024U; // 与 kernel X_BLOCK_ALIGN_NUM 对齐
61- uint64_t xNums = tilingData.quantReduceScatterTilingInfo.bs * tilingData.quantReduceScatterTilingInfo.hiddenSize;
62- uint64_t aivNum = tilingData.quantReduceScatterTilingInfo.aivNum;
63- uint64_t rankSize = static_cast<uint64_t>(runInfo.rankSize);
64- uint64_t xSliceSizeNums = xNums / rankSize;
65- uint64_t perCoreElem = (xSliceSizeNums + aivNum - 1U) / aivNum;
66- uint64_t xPerBlock = (perCoreElem + TARGET_ITER - 1U) / TARGET_ITER;
67- xPerBlock = std::max<uint64_t>(xPerBlock, MIN_BLOCK);
68- xPerBlock = (xPerBlock / ALIGN_BLOCK) * ALIGN_BLOCK;
69- tilingData.quantReduceScatterTilingInfo.xPerBlock = static_cast<uint32_t>(xPerBlock);
70- tilingData.quantReduceScatterTilingInfo.alignBlock = ALIGN_BLOCK;
71-}
72- 
73-static void SetTilingKey(gert::TilingContext *context)
74-{
75- const char *nodeName = context->GetNodeName();
76- // 设置tilingKey模板参数
77- const uint64_t tilingKey = GET_TPL_TILING_KEY(MTE_COMM);
78- context->SetTilingKey(tilingKey);
79- OP_LOGD(nodeName, "tilingKey is [%lu] in quant_reduce_scatter_v2.", tilingKey);
80-}
81- 
82-/**
83- * @brief 校验attr context
84- */
85-static ge::graphStatus CheckMc2Context(gert::TilingContext *context, const char *nodeName,
86- const QuantReduceScatterConfig &config)
87-{
88- const gert::StorageShape *ctxStorageShape = context->GetInputShape(config.CONTEXT_INDEX);
89- OP_TILING_CHECK(ctxStorageShape == nullptr,
90- OP_LOGE(nodeName, "The context shape is null."),
91- return ge::GRAPH_FAILED);
92- 
93- OP_TILING_CHECK(ctxStorageShape->GetStorageShape().GetDimNum() != 1,
94- OP_LOGE(nodeName,
95- "The context shape dim must be 1, but current actual value is: %lu.",
96- ctxStorageShape->GetStorageShape().GetDimNum()),
97- return ge::GRAPH_FAILED);
98- int64_t ctxDim0 = ctxStorageShape->GetStorageShape().GetDim(0);
99- OP_LOGD(nodeName, "The context dim0 is: %ld.", ctxDim0);
100- 
101- auto ctxDesc = context->GetInputDesc(config.CONTEXT_INDEX);
102- OP_TILING_CHECK(ctxDesc == nullptr,
103- OP_LOGE(nodeName, "The context desc is null."),
104- return ge::GRAPH_FAILED);
105- OP_TILING_CHECK(ctxDesc->GetDataType() != ge::DT_INT32,
106- OP_LOGE(nodeName,
107- "The context dataType is invalid, dataType should be int32, but actual value is: %s.",
108- Ops::Base::ToString(ctxDesc->GetDataType()).c_str()),
109- return ge::GRAPH_FAILED);
110- 
111- OP_TILING_CHECK(static_cast<ge::Format>(ge::GetPrimaryFormat(ctxDesc->GetStorageFormat())) != ge::FORMAT_ND,
112- OP_LOGE(nodeName, "The context format is invalid."),
113- return ge::GRAPH_FAILED);
114- 
115- return ge::GRAPH_SUCCESS;
116-}
117- 
118-/**
119- * @brief quant_reduce_scatter_v2算子的tiling函数
120- * @param context: 框架根据input，output，attrs等信息生成tiling需要的context
121- * @return
122- */
123-static ge::graphStatus QuantReduceScatterV2TilingFunc(gert::TilingContext *context)
124-{
125- OP_LOGD("quant_reduce_scatter_v2", "Enter QuantReduceScatterV2TilingFunc.");
126- 
127- OP_TILING_CHECK(context == nullptr,
128- OP_LOGE("quant_reduce_scatter_v2", "failed to get tiling context in quant_reduce_scatter_v2."),
129- return ge::GRAPH_FAILED);
130- const char *nodeName = context->GetNodeName();
131- OP_TILING_CHECK(nodeName == nullptr,
132- OP_LOGE("quant_reduce_scatter_v2", "failed to get nodeName in quant_reduce_scatter_v2."),
133- return ge::GRAPH_FAILED);
134- 
135- QuantReduceScatterTilingData *tilingData = context->GetTilingData<QuantReduceScatterTilingData>();
136- OP_TILING_CHECK(tilingData == nullptr, OP_LOGE(nodeName, "tilingData is nullptr in quant_reduce_scatter_v2."),
137- return ge::GRAPH_FAILED);
138- 
139- TilingRunInfo runInfo = {};
140- OP_TILING_CHECK(QuantReduceScatterUtilTiling::CheckNpuArch(context) != ge::GRAPH_SUCCESS,
141- OP_LOGE(nodeName, "NpuArch is invalid in quant_reduce_scatter_v2."), return ge::GRAPH_FAILED);
142- 
143- QuantReduceScatterConfig config;
144- config.X_INDEX = 1;
145- config.SCALES_INDEX = 2;
146- config.isMc2Context = true;
147- OP_TILING_CHECK(CheckMc2Context(context, nodeName, config) != ge::GRAPH_SUCCESS,
148- OP_LOGE(nodeName, "context check failed in quant_reduce_scatter_v2."), return ge::GRAPH_FAILED);
149- OP_TILING_CHECK(QuantReduceScatterUtilTiling::CheckTilingFunc(context, runInfo,
150- OpType::OP_QUANT_REDUCE_SCATTER, config) != ge::GRAPH_SUCCESS,
151- OP_LOGE(nodeName, "tiling check failed in quant_reduce_scatter_v2."), return ge::GRAPH_FAILED);
152- 
153- SetTilingData(context, *tilingData, config);
154- SetXPerBlock(*tilingData, runInfo);
155- SetTilingKey(context);
156- return ge::GRAPH_SUCCESS;
157-}
158- 
159-struct QuantReduceScatterV2CompileInfo {};
160-
161-ge::graphStatus TilingParseForQuantReduceScatterV2(gert::TilingParseContext *context)
162-{
163- (void)context;
164- return ge::GRAPH_SUCCESS;
165-}
166- 
167-IMPL_OP_OPTILING(QuantReduceScatterV2)
168- .Tiling(QuantReduceScatterV2TilingFunc)
169- .TilingParse<QuantReduceScatterV2CompileInfo>(TilingParseForQuantReduceScatterV2);
170- 
171-} // namespace MC2Tiling
Dmc2/quant_reduce_scatter_v2/op_host/quant_reduce_scatter_v2_def.cpp+0-94
@@ -1,94 +0,0 @@
1-/**
2- * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4- * CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9- */
10- 
11-/*!
12- * \file quant_reduce_scatter_v2_def.cpp
13- * \brief 算子信息库定义
14- */
15-#include <register/op_def_registry.h>
16- 
17-namespace ops {
18-class QuantReduceScatterV2 : public OpDef {
19-public:
20- explicit QuantReduceScatterV2(const char *name) : OpDef(name)
21- {
22- this->Input("context")
23- .ParamType(REQUIRED)
24- .DataTypeList({ge::DT_INT32})
25- .FormatList({ge::FORMAT_ND})
26- .AutoContiguous();
27- this->Input("x")
28- .ParamType(REQUIRED)
29- .DataType({ge::DT_INT8, ge::DT_INT8, ge::DT_INT8,
30- ge::DT_HIFLOAT8, ge::DT_HIFLOAT8, ge::DT_HIFLOAT8,
31- ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E5M2,
32- ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN,
33- ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E5M2,
34- ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN,
35- ge::DT_FLOAT4_E2M1, ge::DT_FLOAT4_E2M1, ge::DT_FLOAT4_E2M1,
36- ge::DT_FLOAT4_E1M2, ge::DT_FLOAT4_E1M2, ge::DT_FLOAT4_E1M2,
37- ge::DT_FLOAT4_E2M1, ge::DT_FLOAT4_E2M1, ge::DT_FLOAT4_E2M1,
38- ge::DT_FLOAT4_E1M2, ge::DT_FLOAT4_E1M2, ge::DT_FLOAT4_E1M2})
39- .FormatList({ge::FORMAT_ND})
40- .AutoContiguous();
41- this->Input("scales")
42- .ParamType(REQUIRED)
43- .DataType({ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT,
44- ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT,
45- ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT,
46- ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT,
47- ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,
48- ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,
49- ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT,
50- ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT,
51- ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0,
52- ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0, ge::DT_FLOAT8_E8M0})
53- .FormatList({ge::FORMAT_ND})
54- .AutoContiguous();
55- 
56- this->Output("out_put")
57- .ParamType(REQUIRED)
58- .DataType({ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT,
59- ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT,
60- ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT,
61- ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT,
62- ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT,
63- ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT,
64- ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT,
65- ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT,
66- ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT,
67- ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT})
68- .FormatList({ge::FORMAT_ND});
69- 
70- this->Attr("hccl_buffer_size").AttrType(REQUIRED).Int();
71- this->Attr("reduce_op").AttrType(OPTIONAL).String("sum");
72- this->Attr("output_dtype")
73- .AttrType(OPTIONAL)
74- .Int(static_cast<int64_t>(ge::DT_BF16)); // 默认值为bf16，check一下对应的枚举值
75- this->Attr("world_size").AttrType(REQUIRED).Int();
76- 
77- // ascend950 AI处理器定义OpAICoreConfig变量，定制化配置参数
78- OpAICoreConfig aicore_config_950;
79- aicore_config_950.DynamicCompileStaticFlag(true)
80- .DynamicFormatFlag(true)
81- .DynamicRankSupportFlag(true)
82- .DynamicShapeSupportFlag(true)
83- .NeedCheckSupportFlag(false)
84- .PrecisionReduceFlag(true)
85- .ExtendCfgInfo("aclnnSupport.value", "support_aclnn")
86- .ExtendCfgInfo("jitCompile.flag", "static_false") // 动态shape，复用二进制，后续图支持后修改
87- .ExtendCfgInfo("multiKernelSupportDynamicGraph.value", "multi_kernel");
88- this->AICore().AddConfig("ascend950", aicore_config_950);
89- }
90-};
91- 
92-OP_ADD(QuantReduceScatterV2);
93- 
94-} // namespace ops
Dmc2/quant_reduce_scatter_v2/op_kernel/quant_reduce_scatter_v2.cpp+0-51
@@ -1,51 +0,0 @@
1-/**
2- * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4- * CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9-**/
10- 
11-/*!
12- * \file quant_reduce_scatter_v2.cpp
13- * \brief
14- */
15- 
16-#if ASC_DEVKIT_MAJOR >= 9
17-#include "basic_api/kernel_basic_intf.h"
18-#else
19-#include "kernel_operator.h"
20-#endif
21- 
22-#if __has_include("../quant_reduce_scatter/quant_reduce_scatter_tiling_data.h")
23-#include "../quant_reduce_scatter/quant_reduce_scatter_tiling_data.h"
24-#include "../quant_reduce_scatter/quant_reduce_scatter_tiling_key.h"
25-#include "../quant_reduce_scatter/quant_reduce_scatter_mte.h"
26-#else
27-#include "../../quant_reduce_scatter/quant_reduce_scatter_tiling_data.h"
28-#include "../../quant_reduce_scatter/quant_reduce_scatter_tiling_key.h"
29-#include "../../quant_reduce_scatter/quant_reduce_scatter_mte.h"
30-#endif
31- 
32-using namespace AscendC;
33-using namespace QuantReduceScatterImpl;
34- 
35-#if defined(__NPU_ARCH__) && (__NPU_ARCH__ == 3510)
36-#endif
37- 
38-template<uint32_t quantReduceScatterCommMode>
39-__global__ __aicore__ void quant_reduce_scatter_v2(GM_ADDR mc2Context, GM_ADDR x, GM_ADDR scales, GM_ADDR output,
40- GM_ADDR workspaceGM, GM_ADDR tilingGM)
41-{
42- KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
43- REGISTER_TILING_DEFAULT(QuantReduceScatterTilingData);
44- GET_TILING_DATA_WITH_STRUCT(QuantReduceScatterTilingData, tilingData, tilingGM);
45- TPipe pipe;
46- if constexpr (quantReduceScatterCommMode == MTE_COMM) {
47- QuantReduceScatterMte<DTYPE_X, DTYPE_SCALES, DTYPE_OUT_PUT> op;
48- op.Init(mc2Context, x, scales, output, &pipe, &tilingData);
49- op.Process();
50- }
51-}
Dmc2/quant_reduce_scatter_v2/tests/CMakeLists.txt+0-16
@@ -1,16 +0,0 @@
1-# -----------------------------------------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-# CANN Open Software License Agreement Version 2.0 (the "License").
5-# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-# See LICENSE in the root of the software repository for the full text of the License.
9-# -----------------------------------------------------------------------------------------------------------
10- 
11-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12-foreach(SUB_DIR ${CURRENT_DIRS})
13- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14- add_subdirectory(${SUB_DIR})
15- endif()
16-endforeach()
Dmc2/quant_reduce_scatter_v2/tests/ut/CMakeLists.txt+0-16
@@ -1,16 +0,0 @@
1-# -----------------------------------------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-# CANN Open Software License Agreement Version 2.0 (the "License").
5-# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-# See LICENSE in the root of the software repository for the full text of the License.
9-# -----------------------------------------------------------------------------------------------------------
10- 
11-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12-foreach(SUB_DIR ${CURRENT_DIRS})
13- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14- add_subdirectory(${SUB_DIR})
15- endif()
16-endforeach()
Dmc2/quant_reduce_scatter_v2/tests/ut/op_api/CMakeLists.txt+0-13
@@ -1,13 +0,0 @@
1-# -----------------------------------------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-# CANN Open Software License Agreement Version 2.0 (the "License").
5-# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-# See LICENSE in the root of the software repository for the full text of the License.
9-# -----------------------------------------------------------------------------------------------------------
10- 
11-if(UT_TEST_ALL OR OP_API_UT)
12- add_modules_ut_sources(UT_NAME ${OP_API_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
13-endif()
Dmc2/quant_reduce_scatter_v2/tests/ut/op_api/test_aclnn_quant_reduce_scatter_v2.cpp+0-9
@@ -1,9 +0,0 @@
1-/**
2- * Copyright (c) Huawei Technologies Co., Ltd. 2026. All rights reserved.
3- * This file is a part of the CANN Open Software.
4- * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9-**/
Dmc2/quant_reduce_scatter_v2/tests/ut/op_host/CMakeLists.txt+0-22
@@ -1,22 +0,0 @@
1-# -----------------------------------------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-# CANN Open Software License Agreement Version 2.0 (the "License").
5-# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-# See LICENSE in the root of the software repository for the full text of the License.
9-# -----------------------------------------------------------------------------------------------------------
10- 
11-if(UT_TEST_ALL OR OP_HOST_UT)
12- if(UT_INFERSHAPE_FLAG)
13- add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14- endif()
15-endif()
16- 
17-file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
18-foreach(SUB_DIR ${CURRENT_DIRS})
19- if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
20- add_subdirectory(${SUB_DIR})
21- endif()
22-endforeach()
Dmc2/quant_reduce_scatter_v2/tests/ut/op_host/arch35/CMakeLists.txt+0-13
@@ -1,13 +0,0 @@
1-# -----------------------------------------------------------------------------------------------------------
2-# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4-# CANN Open Software License Agreement Version 2.0 (the "License").
5-# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8-# See LICENSE in the root of the software repository for the full text of the License.
9-# -----------------------------------------------------------------------------------------------------------
10- 
11-if(UT_TEST_ALL OR OP_HOST_UT)
12- add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
13-endif()
Dmc2/quant_reduce_scatter_v2/tests/ut/op_host/arch35/test_quant_reduce_scatter_v2_tiling.cpp+0-9
@@ -1,9 +0,0 @@
1-/**
2- * Copyright (c) Huawei Technologies Co., Ltd. 2026. All rights reserved.
3- * This file is a part of the CANN Open Software.
4- * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9-**/
Mtests/test_config.yaml+0-19
@@ -1429,25 +1429,6 @@ mc2:
1429 examples: False1429 examples: False
1430 options:1430 options:
1431 - quant_reduce_scatter1431 - quant_reduce_scatter
1432- - quant_reduce_scatter_v2
1433- 
1434- quant_reduce_scatter_v2:
1435- module: True
1436- src:
1437- - mc2/quant_reduce_scatter
1438- - mc2/quant_reduce_scatter_v2
1439- - mc2/common
1440- - mc2/3rd
1441- exclude:
1442- - mc2/quant_reduce_scatter_v2/docs
1443- - mc2/quant_reduce_scatter_v2/README.md
1444- ut_cov_exclude:
1445- - mc2/quant_reduce_scatter_v2/op_graph
1446- - mc2/quant_reduce_scatter_v2/op_kernel
1447- test:
1448- examples: False
1449- options:
1450- - quant_reduce_scatter_v2
1451 1432 
1452 moe_distribute:1433 moe_distribute:
1453 distribute_barrier:1434 distribute_barrier: