modify quant_reduce_scatter hccl_context-ops-transformer-AtomGit

Mbuild.sh+4-0

	@@ -1494,6 +1494,10 @@ if [ -n "${ascend_op_name}" ];then
1494	if [[ "${ascend_op_name}" == "distribute_barrier" ]] && [[ "${ascend_op_name}" != "distribute_barrier_extend" ]]; then	1494	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	fi	1496	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
1497	CUSTOM_OPTION="${CUSTOM_OPTION} -DASCEND_OP_NAME=${ascend_op_name}"	1501	CUSTOM_OPTION="${CUSTOM_OPTION} -DASCEND_OP_NAME=${ascend_op_name}"
1498	if [[ "${ascend_op_name}" != "fused_infer_attention_score" ]] && [[ "${ascend_op_name}" != "incre_flash_attention" ]]; then	1502	if [[ "${ascend_op_name}" != "fused_infer_attention_score" ]] && [[ "${ascend_op_name}" != "incre_flash_attention" ]]; then
1499	CUSTOM_OPTION="${CUSTOM_OPTION} -DENABLE_TILING_SINK=OFF"	1503	CUSTOM_OPTION="${CUSTOM_OPTION} -DENABLE_TILING_SINK=OFF"

Mclassify_rule.yaml+5-0

	@@ -639,6 +639,11 @@ 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.txt	640	- ops/ops-transformer/mc2/quant_reduce_scatter/CMakeLists.txt
641	- ops/ops-transformer/mc2/quant_reduce_scatter/README.md	641	- 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
642	- ops/ops-transformer/mc2/3rd/common/	647	- ops/ops-transformer/mc2/3rd/common/
643	- ops/ops-transformer/mc2/3rd/ops_legacy/	648	- ops/ops-transformer/mc2/3rd/ops_legacy/
644	- ops/ops-transformer/mc2/3rd/CMakeLists.txt	649	- ops/ops-transformer/mc2/3rd/CMakeLists.txt

Mcmake/func.cmake+1-0

	@@ -452,6 +452,7 @@ 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;"
455	"allto_all_matmul;"	456	"allto_all_matmul;"
456	"matmul_allto_all;"	457	"matmul_allto_all;"
457	"attention_to_ffn;"	458	"attention_to_ffn;"

Mmc2/common/op_api/mc2_context.cpp+188-0

	@@ -560,6 +560,194 @@ 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	+
563	// 模板函数显式实例化	751	// 模板函数显式实例化
564	template void Mc2Context::GetHcclLibFunc<void >(void *handle, const std::string &funcName);	752	template void Mc2Context::GetHcclLibFunc<void >(void *handle, const std::string &funcName);
565		753

Mmc2/common/op_api/mc2_context.h+12-0

	@@ -32,6 +32,7 @@
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"
35		36
36	namespace Mc2Aclnn {	37	namespace Mc2Aclnn {
37		38
	@@ -40,6 +41,8 @@ public:
40	static aclnnStatus GetMc2ContextTensor(const char groupEp, const char opName, uint64_t &hcclBuffSize,	41	static aclnnStatus GetMc2ContextTensor(const char groupEp, const char opName, uint64_t &hcclBuffSize,
41	aclTensor *&mc2Context);	42	aclTensor *&mc2Context);
42	static aclnnStatus GetMc2RankSize(const char *groupEp, uint32_t &rankSize);	43	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);
43		46
44	private:	47	private:
45	explicit Mc2Context();	48	explicit Mc2Context();
	@@ -69,6 +72,15 @@ private:
69	aclnnStatus CheckLinks(uint32_t &netLinkNum, CommLink *linksList);	72	aclnnStatus CheckLinks(uint32_t &netLinkNum, CommLink *linksList);
70	aclnnStatus CheckContextCache(const HcclComm &hcclHandle, const std::string &mc2ContextTag,	73	aclnnStatus CheckContextCache(const HcclComm &hcclHandle, const std::string &mc2ContextTag,
71	const CommEngine &engine, void *&ctx, uint64_t &hcclBuffSize);	74	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	+
72	const std::string GetLibPath();	84	const std::string GetLibPath();
73	template <typename T>	85	template <typename T>
74	T GetHcclLibFunc(void *handle, const std::string &funcName);	86	T GetHcclLibFunc(void *handle, const std::string &funcName);

Amc2/common/op_kernel/mc2_quant_reduce_scatter_context.h+36-0

@@ -0,0 +1,36 @@
	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 PARENT_SCOPE)	21	+ set(quant_all_reduce_depends mc2/common mc2/3rd mc2/quant_reduce_scatter mc2/quant_reduce_scatter_v2 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+20-12

Mmc2/quant_all_reduce/op_kernel/quant_all_reduce_tiling_data.h+1-0

	@@ -27,6 +27,7 @@ struct QuantAllReduceTilingInfo {
27	uint64_t totalWinSize; // Win区总大小，即HCCL_BUFFER_SIZE	27	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;
30	};	31	};
31		32
32	struct QuantAllReduceTilingData {	33	struct QuantAllReduceTilingData {

Mmc2/quant_reduce_scatter/op_api/aclnn_quant_reduce_scatter.cpp+12-135

	@@ -13,6 +13,7 @@
13	* \brief	13	* \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"
16	#include "securec.h"	17	#include "securec.h"
17	#include "acl/acl.h"	18	#include "acl/acl.h"
18	#include "common/utils/op_mc2.h"	19	#include "common/utils/op_mc2.h"
	@@ -29,143 +30,19 @@
29		30
30	using namespace op;	31	using namespace op;
31		32
32	-namespace {	33	+extern "C" aclnnStatus aclnnQuantReduceScatterGetWorkspaceSize(const aclTensor* x, const aclTensor* scales,
33	-enum class NnopbaseHcclServerType : uint32_t {	34	+ const char* group, const char* reduceOp,
34	- NNOPBASE_HCCL_SERVER_TYPE_AICPU = 0,	35	+ aclTensor* output,
35	- NNOPBASE_HCCL_SERVER_TYPE_MTE,	36	+ uint64_t* workspaceSize, aclOpExecutor** executor)
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)
65	{	37	{
66	- OP_CHECK_NULL(x, return false);	38	+ OP_LOGD("aclnnQuantReduceScatterGetWorkspaceSize start");
67	- OP_CHECK_NULL(scales, return false);	39	+ return aclnnQuantReduceScatterBaseGetWorkspaceSize(x, scales, group, reduceOp, output,
68	- OP_CHECK_NULL(output, return false);	40	+ workspaceSize, executor);
69	- return true;
70	}	41	}
71		42
72	-// 检查x、scales、output的数据类型是否在算子的支持列表之内	43	+extern "C" aclnnStatus aclnnQuantReduceScatter(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
73	-static bool CheckKGAllDtypesValid(const aclTensor* x, const aclTensor* scales, const aclTensor* output)	44	+ const aclrtStream stream)
74	{	45	{
75	- if (CheckType(x->GetDataType(), X_DTYPE_KG_SUPPORT_LIST) && CheckType(scales->GetDataType(), SCALES_DTYPE_KG_SUPPORT_LIST) &&	46	+ OP_LOGD("aclnnQuantReduceScatter start");
76	- CheckType(output->GetDataType(), OUTPUT_DTYPE_SUPPORT_LIST)) {	47	+ return aclnnQuantReduceScatterBase(workspace, workspaceSize, executor, stream);
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;
171	}	48	}

Amc2/quant_reduce_scatter/op_api/aclnn_quant_reduce_scatter_base.cpp+225-0

@@ -0,0 +1,225 @@
	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	+}

Amc2/quant_reduce_scatter/op_api/aclnn_quant_reduce_scatter_base.h+40-0

@@ -0,0 +1,40 @@
	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+107-76

	@@ -34,18 +34,21 @@ static bool IsContains(const std::vector<uint32_t> &list, uint32_t value)
34	* @param context: 框架根据input，output，attrs等信息生成tiling需要的context	34	* @param context: 框架根据input，output，attrs等信息生成tiling需要的context
35	* @return	35	* @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)
38	{	39	{
39	const char *nodeName = context->GetNodeName();	40	const char *nodeName = context->GetNodeName();
40	const gert::RuntimeAttrs *attrs = context->GetAttrs();	41	const gert::RuntimeAttrs *attrs = context->GetAttrs();
41	OP_TILING_CHECK(attrs == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "attrs"), return ge::GRAPH_FAILED);	42	OP_TILING_CHECK(attrs == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "attrs"), return ge::GRAPH_FAILED);
42	- // 校验group是否为空	43	+ if (!config.isMc2Context) {
43	- const char *groupPtr = attrs->GetAttrPointer<char>(GROUP_INDEX);	44	+ // 校验group是否为空
44	- OP_TILING_CHECK(groupPtr == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "group"), return ge::GRAPH_FAILED);	45	+ const char *groupPtr = attrs->GetAttrPointer<char>(GROUP_INDEX);
45	- OP_TILING_CHECK(std::string(groupPtr).empty(), OP_LOGE_WITH_INVALID_INPUT(nodeName, "group"),	46	+ OP_TILING_CHECK(groupPtr == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "group"), return ge::GRAPH_FAILED);
46	- return ge::GRAPH_FAILED);	47	+ OP_TILING_CHECK(std::string(groupPtr).empty(), OP_LOGE_WITH_INVALID_INPUT(nodeName, "group"),
47	- runInfo.groupPtr = groupPtr;	48	+ return ge::GRAPH_FAILED);
48	- runInfo.group = std::string(groupPtr);	49	+ runInfo.groupPtr = groupPtr;
		50	+ runInfo.group = std::string(groupPtr);
		51	+ }
49	// 校验reduce_op的类型是否为sum	52	// 校验reduce_op的类型是否为sum
50	const char *reduceOpPtr = attrs->GetAttrPointer<char>(REDUCE_OP_INDEX);	53	const char *reduceOpPtr = attrs->GetAttrPointer<char>(REDUCE_OP_INDEX);
51	OP_TILING_CHECK(reduceOpPtr == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "reduce_op"), return ge::GRAPH_FAILED);	54	OP_TILING_CHECK(reduceOpPtr == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "reduce_op"), return ge::GRAPH_FAILED);
	@@ -69,19 +72,27 @@ static ge::graphStatus CheckAttrsInfo(const gert::TilingContext *context, Tiling
69	* @param runInfo: 封装的doTiling所需要的参数	72	* @param runInfo: 封装的doTiling所需要的参数
70	* @return	73	* @return
71	*/	74	*/
72	-static ge::graphStatus SetRankSize(const gert::TilingContext *context, TilingRunInfo &runInfo)	75	+static ge::graphStatus SetRankSize(const gert::TilingContext *context, TilingRunInfo &runInfo,
		76	+ const QuantReduceScatterConfig& config)
73	{	77	{
74	const char *nodeName = context->GetNodeName();	78	const char *nodeName = context->GetNodeName();
75	// attrs在函数CheckAttrsInfo中已做校验	79	// attrs在函数CheckAttrsInfo中已做校验
76	const gert::RuntimeAttrs *attrs = context->GetAttrs();	80	const gert::RuntimeAttrs *attrs = context->GetAttrs();
77	const int64_t *rankSizePtr = attrs->GetAttrPointer<int64_t>(WORLD_SIZE_INDEX);	81	const int64_t *rankSizePtr = attrs->GetAttrPointer<int64_t>(WORLD_SIZE_INDEX);
78	- if (rankSizePtr == nullptr \|\| *rankSizePtr == RANK_SIZE_DEFAULT) {	82	+ if (!config.isMc2Context) {
79	- int64_t rankSize = 0;	83	+ if (rankSizePtr == nullptr \|\| *rankSizePtr == RANK_SIZE_DEFAULT) {
80	- OP_TILING_CHECK(!mc2tiling::GetRankSize(nodeName, runInfo.groupPtr, rankSize),	84	+ int64_t rankSize = 0;
81	- OP_LOGE(nodeName, "Get rankSize failed."),	85	+ OP_TILING_CHECK(!mc2tiling::GetRankSize(nodeName, runInfo.groupPtr, rankSize),
82	- return ge::GRAPH_FAILED);	86	+ OP_LOGE(nodeName, "Get rankSize failed."),
83	- runInfo.rankSize = rankSize;	87	+ return ge::GRAPH_FAILED);
		88	+ runInfo.rankSize = rankSize;
		89	+ } else {
		90	+ runInfo.rankSize = *rankSizePtr;
		91	+ }
84	} else {	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);
85	runInfo.rankSize = *rankSizePtr;	96	runInfo.rankSize = *rankSizePtr;
86	}	97	}
87	OP_TILING_CHECK(std::find(RANK_SIZE_LIST.begin(), RANK_SIZE_LIST.end(), runInfo.rankSize) >= RANK_SIZE_LIST.end(),	98	OP_TILING_CHECK(std::find(RANK_SIZE_LIST.begin(), RANK_SIZE_LIST.end(), runInfo.rankSize) >= RANK_SIZE_LIST.end(),
	@@ -96,12 +107,13 @@ static ge::graphStatus SetRankSize(const gert::TilingContext *context, TilingRun
96	* @param runInfo: 封装的doTiling所需要的参数	107	* @param runInfo: 封装的doTiling所需要的参数
97	* @return	108	* @return
98	*/	109	*/
99	-static bool SetQuantMode(const gert::TilingContext *context, TilingRunInfo &runInfo)	110	+static bool SetQuantMode(const gert::TilingContext *context, TilingRunInfo &runInfo,
		111	+ const QuantReduceScatterConfig& config)
100	{	112	{
101	const char *nodeName = context->GetNodeName();	113	const char *nodeName = context->GetNodeName();
102	// context->GetInputDesc在函数CheckTensorDataType中已经校验	114	// context->GetInputDesc在函数CheckTensorDataType中已经校验
103	- ge::DataType xDtype = context->GetInputDesc(X_INDEX)->GetDataType();	115	+ ge::DataType xDtype = context->GetInputDesc(config.X_INDEX)->GetDataType();
104	- ge::DataType scalesDtype = context->GetInputDesc(SCALES_INDEX)->GetDataType();	116	+ ge::DataType scalesDtype = context->GetInputDesc(config.SCALES_INDEX)->GetDataType();
105	// 0: 无量化模式; 1: TG量化; 2: MX量化	117	// 0: 无量化模式; 1: TG量化; 2: MX量化
106	uint32_t quantMode = 0;	118	uint32_t quantMode = 0;
107	if (IsContains(X_DTYPE_LIST, xDtype) && scalesDtype == ge::DT_FLOAT) {	119	if (IsContains(X_DTYPE_LIST, xDtype) && scalesDtype == ge::DT_FLOAT) {
	@@ -125,20 +137,21 @@ static bool SetQuantMode(const gert::TilingContext *context, TilingRunInfo &runI
125	* @param runInfo: 封装的doTiling所需要的参数	137	* @param runInfo: 封装的doTiling所需要的参数
126	* @return	138	* @return
127	*/	139	*/
128	-static bool CheckTensorDataType(const gert::TilingContext *context, TilingRunInfo &runInfo)	140	+static bool CheckTensorDataType(const gert::TilingContext *context, TilingRunInfo &runInfo,
		L L liangfuzhan6月4日缺少对DT_FLOAT4_E2M1的校验 yifuxiong 6月5日评论： 6月5日评论：
		141	+ const QuantReduceScatterConfig& config)
129	{	142	{
130	const char *nodeName = context->GetNodeName();	143	const char *nodeName = context->GetNodeName();
131	// 校验x的dtype	144	// 校验x的dtype
132	- auto xDesc = context->GetInputDesc(X_INDEX);	145	+ auto xDesc = context->GetInputDesc(config.X_INDEX);
133	OP_TILING_CHECK(xDesc == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "x"), return false);	146	OP_TILING_CHECK(xDesc == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "x"), return false);
134	- ge::DataType xDtype = context->GetInputDesc(X_INDEX)->GetDataType();	147	+ ge::DataType xDtype = context->GetInputDesc(config.X_INDEX)->GetDataType();
135	OP_TILING_CHECK(!IsContains(X_DTYPE_LIST, xDtype),	148	OP_TILING_CHECK(!IsContains(X_DTYPE_LIST, xDtype),
136	OP_LOGE_FOR_INVALID_DTYPE(nodeName, "x", Ops::Base::ToString(xDtype).c_str(), "int8/hifloat8/float8_e4m3fn/float8_e5m2"),	149	OP_LOGE_FOR_INVALID_DTYPE(nodeName, "x", Ops::Base::ToString(xDtype).c_str(), "int8/hifloat8/float8_e4m3fn/float8_e5m2"),
137	return false);	150	return false);
138	// 校验scales的dtype	151	// 校验scales的dtype
139	- auto scalesDesc = context->GetInputDesc(SCALES_INDEX);	152	+ auto scalesDesc = context->GetInputDesc(config.SCALES_INDEX);
140	OP_TILING_CHECK(scalesDesc == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "scales"), return false);	153	OP_TILING_CHECK(scalesDesc == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "scales"), return false);
141	- ge::DataType scalesDtype = context->GetInputDesc(SCALES_INDEX)->GetDataType();	154	+ ge::DataType scalesDtype = context->GetInputDesc(config.SCALES_INDEX)->GetDataType();
142	OP_TILING_CHECK(!IsContains(SCALES_DTYPE_LIST, scalesDtype),	155	OP_TILING_CHECK(!IsContains(SCALES_DTYPE_LIST, scalesDtype),
143	OP_LOGE_FOR_INVALID_DTYPE(nodeName, "scales", Ops::Base::ToString(scalesDtype).c_str(), "float/float8_e8m0"),	156	OP_LOGE_FOR_INVALID_DTYPE(nodeName, "scales", Ops::Base::ToString(scalesDtype).c_str(), "float/float8_e8m0"),
144	return false);	157	return false);
	@@ -150,7 +163,7 @@ static bool CheckTensorDataType(const gert::TilingContext *context, TilingRunInf
150	OP_LOGE_FOR_INVALID_DTYPE(nodeName, "output", Ops::Base::ToString(outputType).c_str(), "float16/bfloat16/float"),	163	OP_LOGE_FOR_INVALID_DTYPE(nodeName, "output", Ops::Base::ToString(outputType).c_str(), "float16/bfloat16/float"),
151	return false);	164	return false);
152	// 设置量化模式	165	// 设置量化模式
153	- OP_TILING_CHECK(!SetQuantMode(context, runInfo), OP_LOGE(nodeName, "get quantMode error."), return false);	166	+ OP_TILING_CHECK(!SetQuantMode(context, runInfo, config), OP_LOGE(nodeName, "get quantMode error."), return false);
154	return true;	167	return true;
155	}	168	}
156		169
	@@ -160,12 +173,13 @@ static bool CheckTensorDataType(const gert::TilingContext *context, TilingRunInf
160	* @param opType: 当前op类型	173	* @param opType: 当前op类型
161	* @return	174	* @return
162	*/	175	*/
163	-static bool CheckXDimValid(const gert::TilingContext *context, const OpType opType)	176	+static bool CheckXDimValid(const gert::TilingContext *context, const OpType opType,
		177	+ const QuantReduceScatterConfig& config)
164	{	178	{
165	(void)opType; // Reserved for future extension	179	(void)opType; // Reserved for future extension
166	const char *nodeName = context->GetNodeName();	180	const char *nodeName = context->GetNodeName();
167	// context->GetInputShape在函数CheckInputTensorDim中已经校验	181	// context->GetInputShape在函数CheckInputTensorDim中已经校验
168	- size_t xDimNum = context->GetInputShape(X_INDEX)->GetStorageShape().GetDimNum();	182	+ size_t xDimNum = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDimNum();
169	// quant_all_reduce和quant_reduce_scatter算子的x可能是2维或者3维，即x.shape(bs, h)或x.shape(b, s, h)	183	// quant_all_reduce和quant_reduce_scatter算子的x可能是2维或者3维，即x.shape(bs, h)或x.shape(b, s, h)
170	bool inValidDimNum = (xDimNum != TWO_DIMS) && (xDimNum != THREE_DIMS);	184	bool inValidDimNum = (xDimNum != TWO_DIMS) && (xDimNum != THREE_DIMS);
171	OP_TILING_CHECK(inValidDimNum,	185	OP_TILING_CHECK(inValidDimNum,
	@@ -182,12 +196,13 @@ static bool CheckXDimValid(const gert::TilingContext *context, const OpType opTy
182	* @param opType: 当前op类型	196	* @param opType: 当前op类型
183	* @return	197	* @return
184	*/	198	*/
185	-static bool CheckXShapeValid(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType)	199	+static bool CheckXShapeValid(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType,
		200	+ const QuantReduceScatterConfig& config)
186	{	201	{
187	(void)opType; // Reserved for future extension	202	(void)opType; // Reserved for future extension
188	const char *nodeName = context->GetNodeName();	203	const char *nodeName = context->GetNodeName();
189	// context->GetInputShape在函数CheckInputTensorDim中已经校验	204	// context->GetInputShape在函数CheckInputTensorDim中已经校验
190	- const gert::StorageShape *xShape = context->GetInputShape(X_INDEX);	205	+ const gert::StorageShape *xShape = context->GetInputShape(config.X_INDEX);
191	// 获取x各维度值	206	// 获取x各维度值
192	size_t xDimNum = xShape->GetStorageShape().GetDimNum();	207	size_t xDimNum = xShape->GetStorageShape().GetDimNum();
193	uint64_t xValueOne = xShape->GetStorageShape().GetDim(DIM_ZERO);	208	uint64_t xValueOne = xShape->GetStorageShape().GetDim(DIM_ZERO);
	@@ -200,7 +215,7 @@ static bool CheckXShapeValid(const gert::TilingContext *context, TilingRunInfo &
200	// 当x是3维时，x.shape = (B, S, H)	215	// 当x是3维时，x.shape = (B, S, H)
201	if (xDimNum == THREE_DIMS) {	216	if (xDimNum == THREE_DIMS) {
202	xValueBS = xValueOne * xValueTwo;	217	xValueBS = xValueOne * xValueTwo;
203	- xValueH = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_TWO);	218	+ xValueH = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_TWO);
204	emptyTensor = emptyTensor \|\| xValueH == 0;	219	emptyTensor = emptyTensor \|\| xValueH == 0;
205	}	220	}
206		221
	@@ -236,9 +251,10 @@ static bool CheckXShapeValid(const gert::TilingContext *context, TilingRunInfo &
236	* @return 计算出的正确scales维度向量	251	* @return 计算出的正确scales维度向量
237	*/	252	*/
238	static std::vector<uint64_t> CalculateExpectedScalesShape(const gert::TilingContext *context,	253	static std::vector<uint64_t> CalculateExpectedScalesShape(const gert::TilingContext *context,
239	- TilingRunInfo &runInfo)	254	+ TilingRunInfo &runInfo,
		255	+ const QuantReduceScatterConfig& config)
240	{	256	{
241	- const gert::StorageShape *xShape = context->GetInputShape(X_INDEX);	257	+ const gert::StorageShape *xShape = context->GetInputShape(config.X_INDEX);
242		258
243	// 获取x的维度和值	259	// 获取x的维度和值
244	size_t xDimNum = xShape->GetStorageShape().GetDimNum();	260	size_t xDimNum = xShape->GetStorageShape().GetDimNum();
	@@ -305,10 +321,10 @@ static std::string FormatShape(const std::vector<uint64_t> &dims)
305	*/	321	*/
306	static bool CheckScalesValid(const gert::TilingContext *context,	322	static bool CheckScalesValid(const gert::TilingContext *context,
307	const std::vector<uint64_t> &expectedScalesDims,	323	const std::vector<uint64_t> &expectedScalesDims,
308	- const TilingRunInfo &runInfo)	324	+ const TilingRunInfo &runInfo, const QuantReduceScatterConfig& config)
309	{	325	{
310	const char *nodeName = context->GetNodeName();	326	const char *nodeName = context->GetNodeName();
311	- const gert::StorageShape *scalesShape = context->GetInputShape(SCALES_INDEX);	327	+ const gert::StorageShape *scalesShape = context->GetInputShape(config.SCALES_INDEX);
312		328
313	// 将quantMode转换为可读字符串	329	// 将quantMode转换为可读字符串
314	const char* quantModeStr = "";	330	const char* quantModeStr = "";
	@@ -348,26 +364,29 @@ static bool CheckScalesValid(const gert::TilingContext *context,
348	* @param opType:当前op类型	364	* @param opType:当前op类型
349	* @return	365	* @return
350	*/	366	*/
351	-static bool CheckInputTensorDim(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType)	367	+static bool CheckInputTensorDim(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType,
		368	+ const QuantReduceScatterConfig& config)
352	{	369	{
353	const char *nodeName = context->GetNodeName();	370	const char *nodeName = context->GetNodeName();
354	// 1.校验x相关	371	// 1.校验x相关
355	- const gert::StorageShape *xShape = context->GetInputShape(X_INDEX);	372	+ const gert::StorageShape *xShape = context->GetInputShape(config.X_INDEX);
356	// 校验x不为空	373	// 校验x不为空
357	OP_TILING_CHECK(xShape == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "xShape"), return false);	374	OP_TILING_CHECK(xShape == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "xShape"), return false);
358	// 校验x维度数量合法性	375	// 校验x维度数量合法性
359	- OP_TILING_CHECK(!CheckXDimValid(context, opType), OP_LOGE(nodeName, "x dimensions is invalid."), return false);	376	+ OP_TILING_CHECK(!CheckXDimValid(context, opType, config),
		377	+ OP_LOGE(nodeName, "x dimensions is invalid."), return false);
360	// 校验x.shape合法性	378	// 校验x.shape合法性
361	- OP_TILING_CHECK(!CheckXShapeValid(context, runInfo, opType), OP_LOGE(nodeName, "x shapes is invalid."), return false);	379	+ OP_TILING_CHECK(!CheckXShapeValid(context, runInfo, opType, config),
		380	+ OP_LOGE(nodeName, "x shapes is invalid."), return false);
362		381
363	// 2.校验scales	382	// 2.校验scales
364	- const gert::StorageShape *scalesShape = context->GetInputShape(SCALES_INDEX);	383	+ const gert::StorageShape *scalesShape = context->GetInputShape(config.SCALES_INDEX);
365	// 根据x计算正确的scales, 当scale形状不匹配时，会打印预期的形状和实际的形状	384	// 根据x计算正确的scales, 当scale形状不匹配时，会打印预期的形状和实际的形状
366	- std::vector<uint64_t> expectedScalesDims = CalculateExpectedScalesShape(context, runInfo);	385	+ std::vector<uint64_t> expectedScalesDims = CalculateExpectedScalesShape(context, runInfo, config);
367	// 校验scales不为空	386	// 校验scales不为空
368	OP_TILING_CHECK(scalesShape == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "scales"), return false);	387	OP_TILING_CHECK(scalesShape == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "scales"), return false);
369	// 校验scales维度和shape是否正确	388	// 校验scales维度和shape是否正确
370	- OP_TILING_CHECK(!CheckScalesValid(context, expectedScalesDims, runInfo),	389	+ OP_TILING_CHECK(!CheckScalesValid(context, expectedScalesDims, runInfo, config),
371	OP_LOGE(nodeName, "scales dimensions and shapes is invalid in the quantmode."), return false);	390	OP_LOGE(nodeName, "scales dimensions and shapes is invalid in the quantmode."), return false);
372	return true;	391	return true;
373	}	392	}
	@@ -400,13 +419,14 @@ static bool CheckOutputDimSize(const gert::TilingContext *context, size_t output
400	* @brief 检查quant_all_reduce的输出形状	419	* @brief 检查quant_all_reduce的输出形状
401	*/	420	*/
402	static bool CheckAllReduceOutputShape(const gert::TilingContext context, const gert::StorageShape outputShape,	421	static bool CheckAllReduceOutputShape(const gert::TilingContext context, const gert::StorageShape outputShape,
403	- size_t outputDim, size_t xDimNum, TilingRunInfo &runInfo, const char *nodeName)	422	+ size_t outputDim, size_t xDimNum, TilingRunInfo &runInfo, const char *nodeName,
		423	+ const QuantReduceScatterConfig& config)
404	{	424	{
405	(void)xDimNum; // Reserved for future extension	425	(void)xDimNum; // Reserved for future extension
406	uint64_t outputValueOne = outputShape->GetStorageShape().GetDim(DIM_ZERO);	426	uint64_t outputValueOne = outputShape->GetStorageShape().GetDim(DIM_ZERO);
407	uint64_t outputValueTwo = outputShape->GetStorageShape().GetDim(DIM_ONE);	427	uint64_t outputValueTwo = outputShape->GetStorageShape().GetDim(DIM_ONE);
408	- uint64_t xValueOne = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);	428	+ uint64_t xValueOne = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);
409	- uint64_t xValueTwo = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ONE);	429	+ uint64_t xValueTwo = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
410		430
411	// 对于quant_all_reduce算子，output.shape必须等于x.shape	431	// 对于quant_all_reduce算子，output.shape必须等于x.shape
412	bool invalidShape = (xValueOne != outputValueOne) \|\| (xValueTwo != outputValueTwo); // 校验前两维的大小	432	bool invalidShape = (xValueOne != outputValueOne) \|\| (xValueTwo != outputValueTwo); // 校验前两维的大小
	@@ -414,7 +434,7 @@ static bool CheckAllReduceOutputShape(const gert::TilingContext *context, const
414	// quant_all_reduce算子支持三维，output可能需要校验第3维	434	// quant_all_reduce算子支持三维，output可能需要校验第3维
415	if (outputDim == THREE_DIMS) {	435	if (outputDim == THREE_DIMS) {
416	uint64_t outputValueThree = outputShape->GetStorageShape().GetDim(DIM_TWO);	436	uint64_t outputValueThree = outputShape->GetStorageShape().GetDim(DIM_TWO);
417	- uint64_t xValueThree = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_TWO);	437	+ uint64_t xValueThree = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_TWO);
418	OP_LOGI(nodeName, "output dim2 is %lu, x dim2 is %lu", outputValueThree, xValueThree);	438	OP_LOGI(nodeName, "output dim2 is %lu, x dim2 is %lu", outputValueThree, xValueThree);
419	invalidShape = invalidShape \|\| (xValueThree != outputValueThree); // 校验第三维的大小	439	invalidShape = invalidShape \|\| (xValueThree != outputValueThree); // 校验第三维的大小
420	OP_TILING_CHECK(invalidShape,	440	OP_TILING_CHECK(invalidShape,
	@@ -439,11 +459,12 @@ static bool CheckAllReduceOutputShape(const gert::TilingContext *context, const
439	static bool CheckReduceScatter3DShape(const gert::TilingContext *context,	459	static bool CheckReduceScatter3DShape(const gert::TilingContext *context,
440	uint64_t outputValueOne, uint64_t outputValueTwo,	460	uint64_t outputValueOne, uint64_t outputValueTwo,
441	uint64_t xValueOne, uint64_t xValueTwo,	461	uint64_t xValueOne, uint64_t xValueTwo,
442	- TilingRunInfo &runInfo, const char *nodeName)	462	+ TilingRunInfo &runInfo, const char *nodeName,
		463	+ const QuantReduceScatterConfig& config)
443	{	464	{
444	// 若X为3维，则要对b,s进行合轴，再与output判断是否合法	465	// 若X为3维，则要对b,s进行合轴，再与output判断是否合法
445	uint64_t xValueBS = xValueOne * xValueTwo;	466	uint64_t xValueBS = xValueOne * xValueTwo;
446	- uint64_t xValueThree = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_TWO);	467	+ uint64_t xValueThree = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_TWO);
447	bool invalidShape = xValueBS / runInfo.rankSize != outputValueOne; // 校验bs轴	468	bool invalidShape = xValueBS / runInfo.rankSize != outputValueOne; // 校验bs轴
448	invalidShape = invalidShape \|\| (xValueThree != outputValueTwo); // 校验h轴	469	invalidShape = invalidShape \|\| (xValueThree != outputValueTwo); // 校验h轴
449	OP_TILING_CHECK(invalidShape,	470	OP_TILING_CHECK(invalidShape,
	@@ -476,18 +497,19 @@ static bool CheckReduceScatter2DShape(uint64_t outputValueOne, uint64_t outputVa
476	* @brief 检查quant_reduce_scatter的输出形状	497	* @brief 检查quant_reduce_scatter的输出形状
477	*/	498	*/
478	static bool CheckReduceScatterOutputShape(const gert::TilingContext context, const gert::StorageShape outputShape,	499	static bool CheckReduceScatterOutputShape(const gert::TilingContext context, const gert::StorageShape outputShape,
479	- size_t outputDim, size_t xDimNum, TilingRunInfo &runInfo, const char *nodeName)	500	+ size_t outputDim, size_t xDimNum, TilingRunInfo &runInfo,
		501	+ const char *nodeName, const QuantReduceScatterConfig& config)
480	{	502	{
481	(void)outputDim; // Reserved for future extension	503	(void)outputDim; // Reserved for future extension
482	uint64_t outputValueOne = outputShape->GetStorageShape().GetDim(DIM_ZERO);	504	uint64_t outputValueOne = outputShape->GetStorageShape().GetDim(DIM_ZERO);
483	uint64_t outputValueTwo = outputShape->GetStorageShape().GetDim(DIM_ONE);	505	uint64_t outputValueTwo = outputShape->GetStorageShape().GetDim(DIM_ONE);
484	- uint64_t xValueOne = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);	506	+ uint64_t xValueOne = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);
485	- uint64_t xValueTwo = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ONE);	507	+ uint64_t xValueTwo = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
486		508
487	// 对于quant_reduce_scatter算子, 输出output一定为2维，判断x维度大小决定是否b,s合轴	509	// 对于quant_reduce_scatter算子, 输出output一定为2维，判断x维度大小决定是否b,s合轴
488	if (xDimNum == THREE_DIMS) {	510	if (xDimNum == THREE_DIMS) {
489	return CheckReduceScatter3DShape(context, outputValueOne, outputValueTwo,	511	return CheckReduceScatter3DShape(context, outputValueOne, outputValueTwo,
490	- xValueOne, xValueTwo, runInfo, nodeName);	512	+ xValueOne, xValueTwo, runInfo, nodeName, config);
491	} else {	513	} else {
492	return CheckReduceScatter2DShape(outputValueOne, outputValueTwo,	514	return CheckReduceScatter2DShape(outputValueOne, outputValueTwo,
493	xValueOne, xValueTwo, runInfo, nodeName);	515	xValueOne, xValueTwo, runInfo, nodeName);
	@@ -500,14 +522,15 @@ static bool CheckReduceScatterOutputShape(const gert::TilingContext *context, co
500	* @param runInfo: 封装的doTiling所需要的参数	522	* @param runInfo: 封装的doTiling所需要的参数
501	* @return	523	* @return
502	*/	524	*/
503	-static bool CheckOutputDim(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType)	525	+static bool CheckOutputDim(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType,
		526	+ const QuantReduceScatterConfig& config)
504	{	527	{
505	const char *nodeName = context->GetNodeName();	528	const char *nodeName = context->GetNodeName();
506	// context->GetOutputShape在函数CheckOutputTensorDim中已经校验	529	// context->GetOutputShape在函数CheckOutputTensorDim中已经校验
507	const gert::StorageShape *outputShape = context->GetOutputShape(OUTPUT_INDEX);	530	const gert::StorageShape *outputShape = context->GetOutputShape(OUTPUT_INDEX);
508	size_t outputDim = outputShape->GetStorageShape().GetDimNum();	531	size_t outputDim = outputShape->GetStorageShape().GetDimNum();
509	// context->GetInputShape在函数CheckInputTensorDim中已经校验	532	// context->GetInputShape在函数CheckInputTensorDim中已经校验
510	- size_t xDimNum = context->GetInputShape(X_INDEX)->GetStorageShape().GetDimNum();	533	+ size_t xDimNum = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDimNum();
511	// 检查output的维度大小	534	// 检查output的维度大小
512	if (!CheckOutputDimSize(context, outputDim, xDimNum, opType, nodeName)) {	535	if (!CheckOutputDimSize(context, outputDim, xDimNum, opType, nodeName)) {
513	return false;	536	return false;
	@@ -515,9 +538,9 @@ static bool CheckOutputDim(const gert::TilingContext *context, TilingRunInfo &ru
515		538
516	// 检查输出output形状与输入x形状的关系	539	// 检查输出output形状与输入x形状的关系
517	if (opType == OpType::OP_QUANT_ALL_REDUCE) {	540	if (opType == OpType::OP_QUANT_ALL_REDUCE) {
518	- return CheckAllReduceOutputShape(context, outputShape, outputDim, xDimNum, runInfo, nodeName);	541	+ return CheckAllReduceOutputShape(context, outputShape, outputDim, xDimNum, runInfo, nodeName, config);
519	} else {	542	} else {
520	- return CheckReduceScatterOutputShape(context, outputShape, outputDim, xDimNum, runInfo, nodeName);	543	+ return CheckReduceScatterOutputShape(context, outputShape, outputDim, xDimNum, runInfo, nodeName, config);
521	}	544	}
522	}	545	}
523		546
	@@ -528,13 +551,14 @@ static bool CheckOutputDim(const gert::TilingContext *context, TilingRunInfo &ru
528	* @param opType: 当前op类型	551	* @param opType: 当前op类型
529	* @return	552	* @return
530	*/	553	*/
531	-static bool CheckOutputTensorDim(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType)	554	+static bool CheckOutputTensorDim(const gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType,
		555	+ const QuantReduceScatterConfig& config)
532	{	556	{
533	const char *nodeName = context->GetNodeName();	557	const char *nodeName = context->GetNodeName();
534	// 红线校验	558	// 红线校验
535	const gert::StorageShape *outputShape = context->GetOutputShape(OUTPUT_INDEX);	559	const gert::StorageShape *outputShape = context->GetOutputShape(OUTPUT_INDEX);
536	OP_TILING_CHECK(outputShape == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "output"), return false);	560	OP_TILING_CHECK(outputShape == nullptr, OP_LOGE_WITH_INVALID_INPUT(nodeName, "output"), return false);
537	- return CheckOutputDim(context, runInfo, opType);	561	+ return CheckOutputDim(context, runInfo, opType, config);
538	}	562	}
539		563
540	/**	564	/**
	@@ -542,17 +566,17 @@ static bool CheckOutputTensorDim(const gert::TilingContext *context, TilingRunIn
542	* @param context: 框架根据input，output，attrs等信息生成tiling需要的context	566	* @param context: 框架根据input，output，attrs等信息生成tiling需要的context
543	* @return	567	* @return
544	*/	568	*/
545	-static bool CheckTensorFormat(const gert::TilingContext *context)	569	+static bool CheckTensorFormat(const gert::TilingContext *context, const QuantReduceScatterConfig& config)
546	{	570	{
547	const char *nodeName = context->GetNodeName();	571	const char *nodeName = context->GetNodeName();
548	// context->GetInputDesc在CheckTensorDataType函数中已经校验	572	// context->GetInputDesc在CheckTensorDataType函数中已经校验
549	- auto xDesc = context->GetInputDesc(X_INDEX);	573	+ auto xDesc = context->GetInputDesc(config.X_INDEX);
550	ge::Format xFormat = static_cast<ge::Format>(ge::GetPrimaryFormat(xDesc->GetStorageFormat()));	574	ge::Format xFormat = static_cast<ge::Format>(ge::GetPrimaryFormat(xDesc->GetStorageFormat()));
551	OP_TILING_CHECK(	575	OP_TILING_CHECK(
552	xFormat != ge::FORMAT_ND,	576	xFormat != ge::FORMAT_ND,
553	OP_LOGE_FOR_INVALID_FORMAT(nodeName, "x", Ops::Base::ToString(xFormat).c_str(), "ND"),	577	OP_LOGE_FOR_INVALID_FORMAT(nodeName, "x", Ops::Base::ToString(xFormat).c_str(), "ND"),
554	return false);	578	return false);
555	- auto scalesDesc = context->GetInputDesc(SCALES_INDEX);	579	+ auto scalesDesc = context->GetInputDesc(config.SCALES_INDEX);
556	ge::Format scalesFormat = static_cast<ge::Format>(ge::GetPrimaryFormat(scalesDesc->GetStorageFormat()));	580	ge::Format scalesFormat = static_cast<ge::Format>(ge::GetPrimaryFormat(scalesDesc->GetStorageFormat()));
557	OP_TILING_CHECK(scalesFormat != ge::FORMAT_ND,	581	OP_TILING_CHECK(scalesFormat != ge::FORMAT_ND,
558	OP_LOGE_FOR_INVALID_FORMAT(nodeName, "scales", Ops::Base::ToString(scalesFormat).c_str(), "ND"),	582	OP_LOGE_FOR_INVALID_FORMAT(nodeName, "scales", Ops::Base::ToString(scalesFormat).c_str(), "ND"),
	@@ -572,24 +596,25 @@ static bool CheckTensorFormat(const gert::TilingContext *context)
572	* @param runInfo: 封装的doTiling所需要的参数	596	* @param runInfo: 封装的doTiling所需要的参数
573	* @return	597	* @return
574	*/	598	*/
575	-static bool CheckWindowSize(const gert::TilingContext *context, const TilingRunInfo &runInfo)	599	+static bool CheckWindowSize(const gert::TilingContext *context, const TilingRunInfo &runInfo,
		600	+ const QuantReduceScatterConfig& config)
576	{	601	{
577	const char *nodeName = context->GetNodeName();	602	const char *nodeName = context->GetNodeName();
578	// 获取量化模式，数据类型	603	// 获取量化模式，数据类型
579	- uint64_t xValueOne = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);	604	+ uint64_t xValueOne = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);
580	- uint64_t xValueTwo = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ONE);	605	+ uint64_t xValueTwo = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
581	- uint64_t scalesValueOne = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(DIM_ZERO);	606	+ uint64_t scalesValueOne = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_ZERO);
582	- uint64_t scalesValueTwo = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(DIM_ONE);	607	+ uint64_t scalesValueTwo = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_ONE);
583		608
584	// 计算xDataSize	609	// 计算xDataSize
585	uint64_t xValue = xValueOne * xValueTwo;	610	uint64_t xValue = xValueOne * xValueTwo;
586	uint64_t scalesValue = scalesValueOne * scalesValueTwo;	611	uint64_t scalesValue = scalesValueOne * scalesValueTwo;
587	uint32_t scalesLastDim = DIM_TWO;	612	uint32_t scalesLastDim = DIM_TWO;
588	- size_t xDimNum = context->GetInputShape(X_INDEX)->GetStorageShape().GetDimNum();	613	+ size_t xDimNum = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDimNum();
589	if (xDimNum == THREE_DIMS) {	614	if (xDimNum == THREE_DIMS) {
590	- uint64_t xValueThree = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_TWO);	615	+ uint64_t xValueThree = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_TWO);
591	xValue = xValue * xValueThree;	616	xValue = xValue * xValueThree;
592	- uint64_t scalesValueThree = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(DIM_TWO);	617	+ uint64_t scalesValueThree = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_TWO);
593	scalesValue = scalesValue * scalesValueThree;	618	scalesValue = scalesValue * scalesValueThree;
594	scalesLastDim = DIM_THREE;	619	scalesLastDim = DIM_THREE;
595	}	620	}
	@@ -603,7 +628,7 @@ static bool CheckWindowSize(const gert::TilingContext *context, const TilingRunI
603	scalesSize = scalesValue * SCALE_DTYPE_SIZE_FOUR;	628	scalesSize = scalesValue * SCALE_DTYPE_SIZE_FOUR;
604	} else if (runInfo.quantMode == MX_QUANT_MOD) {	629	} else if (runInfo.quantMode == MX_QUANT_MOD) {
605	// scales的最后一维一定为2	630	// scales的最后一维一定为2
606	- uint64_t scalesValueLast = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(scalesLastDim);	631	+ uint64_t scalesValueLast = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(scalesLastDim);
607	scalesSize = scalesValue * scalesValueLast * SCALE_DTYPE_SIZE_ONE;	632	scalesSize = scalesValue * scalesValueLast * SCALE_DTYPE_SIZE_ONE;
608	}	633	}
609	uint64_t scalesDataSize = ((scalesSize + WIN_ADDR_ALIGN - 1UL) / WIN_ADDR_ALIGN) * WIN_ADDR_ALIGN;	634	uint64_t scalesDataSize = ((scalesSize + WIN_ADDR_ALIGN - 1UL) / WIN_ADDR_ALIGN) * WIN_ADDR_ALIGN;
	@@ -663,25 +688,31 @@ ge::graphStatus QuantReduceScatterUtilTiling::CheckNpuArch(const gert::TilingCon
663	* @return	688	* @return
664	*/	689	*/
665	ge::graphStatus QuantReduceScatterUtilTiling::CheckTilingFunc(gert::TilingContext *context, TilingRunInfo &runInfo,	690	ge::graphStatus QuantReduceScatterUtilTiling::CheckTilingFunc(gert::TilingContext *context, TilingRunInfo &runInfo,
666	- const OpType opType)	691	+ const OpType opType,
		692	+ const QuantReduceScatterConfig& config)
667	{	693	{
668	const char *nodeName = context->GetNodeName();	694	const char *nodeName = context->GetNodeName();
669	// set group	695	// set group
670	- OP_TILING_CHECK(CheckAttrsInfo(context, runInfo) != ge::GRAPH_SUCCESS, OP_LOGE(nodeName, "attrs are invalied."),	696	+ OP_TILING_CHECK(CheckAttrsInfo(context, runInfo, config) != ge::GRAPH_SUCCESS,
		697	+ OP_LOGE(nodeName, "attrs are invalied."),
671	return ge::GRAPH_FAILED);	698	return ge::GRAPH_FAILED);
672	// set rankSize	699	// set rankSize
673	- OP_TILING_CHECK(SetRankSize(context, runInfo) != ge::GRAPH_SUCCESS, OP_LOGE(nodeName, "set rankSize failed."),	700	+ OP_TILING_CHECK(SetRankSize(context, runInfo, config) != ge::GRAPH_SUCCESS,
		701	+ OP_LOGE(nodeName, "set rankSize failed."),
674	return ge::GRAPH_FAILED);	702	return ge::GRAPH_FAILED);
675	// set quantMode	703	// set quantMode
676	- OP_TILING_CHECK(!CheckTensorDataType(context, runInfo), OP_LOGE(nodeName, "tensor datatype is invalid."),	704	+ OP_TILING_CHECK(!CheckTensorDataType(context, runInfo, config),
		705	+ OP_LOGE(nodeName, "tensor datatype is invalid."),
677	return ge::GRAPH_FAILED);	706	return ge::GRAPH_FAILED);
678	- OP_TILING_CHECK(!CheckInputTensorDim(context, runInfo, opType), OP_LOGE(nodeName, "input tensor dim is invalid."),	707	+ OP_TILING_CHECK(!CheckInputTensorDim(context, runInfo, opType, config),
		708	+ OP_LOGE(nodeName, "input tensor dim is invalid."),
679	return ge::GRAPH_FAILED);	709	return ge::GRAPH_FAILED);
680	- OP_TILING_CHECK(!CheckOutputTensorDim(context, runInfo, opType), OP_LOGE(nodeName, "output tensor dim is invalid."),	710	+ OP_TILING_CHECK(!CheckOutputTensorDim(context, runInfo, opType, config),
		711	+ OP_LOGE(nodeName, "output tensor dim is invalid."),
681	return ge::GRAPH_FAILED);	712	return ge::GRAPH_FAILED);
682	- OP_TILING_CHECK(!CheckTensorFormat(context), OP_LOGE(nodeName, "tensor format is invalid."),	713	+ OP_TILING_CHECK(!CheckTensorFormat(context, config), OP_LOGE(nodeName, "tensor format is invalid."),
683	return ge::GRAPH_FAILED);	714	return ge::GRAPH_FAILED);
684	- OP_TILING_CHECK(!CheckWindowSize(context, runInfo), OP_LOGE(nodeName, "HCCL_BUFFSIZE is too small."),	715	+ OP_TILING_CHECK(!CheckWindowSize(context, runInfo, config), OP_LOGE(nodeName, "HCCL_BUFFSIZE is too small."),
685	return ge::GRAPH_FAILED);	716	return ge::GRAPH_FAILED);
686	OP_TILING_CHECK(SetWorkSpace(context) != ge::GRAPH_SUCCESS, OP_LOGE(nodeName, "set workspace failed."),	717	OP_TILING_CHECK(SetWorkSpace(context) != ge::GRAPH_SUCCESS, OP_LOGE(nodeName, "set workspace failed."),
687	return ge::GRAPH_FAILED);	718	return ge::GRAPH_FAILED);

Mmc2/quant_reduce_scatter/op_host/op_tiling/common/quant_reduce_scatter_util_tiling.h+8-3

	@@ -21,8 +21,12 @@ using namespace ge;
21	using namespace gert;	21	using namespace gert;
22		22
23	// input index	23	// input index
24	-constexpr size_t X_INDEX = 0;	24	+struct QuantReduceScatterConfig {
25	-constexpr size_t SCALES_INDEX = 1;	25	+ uint64_t CONTEXT_INDEX = 0;
		26	+ uint64_t X_INDEX = 0;
		27	+ uint64_t SCALES_INDEX = 1;
		28	+ bool isMc2Context = false;
		29	+};
26	// output index	30	// output index
27	constexpr size_t OUTPUT_INDEX = 0;	31	constexpr size_t OUTPUT_INDEX = 0;
28	// attr index	32	// attr index
	@@ -84,7 +88,8 @@ struct TilingRunInfo {
84	class QuantReduceScatterUtilTiling {	88	class QuantReduceScatterUtilTiling {
85	public:	89	public:
86	static ge::graphStatus CheckNpuArch(const gert::TilingContext *context);	90	static ge::graphStatus CheckNpuArch(const gert::TilingContext *context);
87	- static ge::graphStatus CheckTilingFunc(gert::TilingContext *context, TilingRunInfo &runInfo, const OpType opType);	91	+ static ge::graphStatus CheckTilingFunc(gert::TilingContext *context, TilingRunInfo &runInfo,
		92	+ const OpType opType, const QuantReduceScatterConfig& config);
88	};	93	};
89		94
90	}; // namespace MC2Tiling	95	}; // namespace MC2Tiling

Mmc2/quant_reduce_scatter/op_host/op_tiling/quant_reduce_scatter_tiling.cpp+18-11

Mmc2/quant_reduce_scatter/op_kernel/mte_comm.h+54-16

	@@ -18,11 +18,15 @@
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"
21	#include "../../common/op_kernel/moe_distribute_base.h"	22	#include "../../common/op_kernel/moe_distribute_base.h"
		23	+#include "../../common/op_kernel/mc2_quant_reduce_scatter_context.h"
22		24
23	namespace QuantMTECommImpl {	25	namespace QuantMTECommImpl {
24		26
25	using namespace AscendC;	27	using namespace AscendC;
		28	+using namespace Mc2Aclnn;
		29	+
26	// 后缀_BYTES 表示单位为字节大小B, _NUM 表示单位为个	30	// 后缀_BYTES 表示单位为字节大小B, _NUM 表示单位为个
27	constexpr static uint32_t UB_ALIGN_BYTES = 32U; // UB按32B对齐	31	constexpr static uint32_t UB_ALIGN_BYTES = 32U; // UB按32B对齐
28	constexpr uint32_t FLOAT_UB_ALIGN_NUM = 8U; // float格式下32B对齐需要 32/4 =8个	32	constexpr uint32_t FLOAT_UB_ALIGN_NUM = 8U; // float格式下32B对齐需要 32/4 =8个
	@@ -41,6 +45,7 @@ template<TemplateTypeClass>
41	class MTECommunication {	45	class MTECommunication {
42	public:	46	public:
43	__aicore__ inline MTECommunication() {};	47	__aicore__ inline MTECommunication() {};
		48	+ __aicore__ inline void InitMc2Context(GM_ADDR mc2Context, const QuantReduceScatterTilingData *tilingData);
44	__aicore__ inline void InitHcclContext();	49	__aicore__ inline void InitHcclContext();
45	__aicore__ inline void InitParams();	50	__aicore__ inline void InitParams();
46	__aicore__ inline void InitGMTensor(GM_ADDR x, GM_ADDR scales, GM_ADDR output, uint64_t alignedXSize, uint64_t dataSpaceGmSize);	51	__aicore__ inline void InitGMTensor(GM_ADDR x, GM_ADDR scales, GM_ADDR output, uint64_t alignedXSize, uint64_t dataSpaceGmSize);
	@@ -57,7 +62,12 @@ public:
57	__aicore__ inline GM_ADDR GetWinDataAddrGm(uint32_t rankId, uint32_t winFlag);	62	__aicore__ inline GM_ADDR GetWinDataAddrGm(uint32_t rankId, uint32_t winFlag);
58	__aicore__ inline GM_ADDR GetWinStatusAddrGm(uint32_t rankId, uint32_t winFlag);	63	__aicore__ inline GM_ADDR GetWinStatusAddrGm(uint32_t rankId, uint32_t winFlag);
59		64
60	- __gm__ Mc2Kernel::HcclOpParam *hcclContext_;	65	+ uint32_t rankIdHccl_{0};
		66	+ uint32_t rankDimHccl_{0};
		67	+ bool isMc2Context_ = false;
		68	+ __gm__ Mc2QuantReduceScatterContext* mc2Context_{nullptr};
		69	+ __gm__ Mc2Kernel::HcclOpParam *hcclContext_{nullptr};
		70	+
61	uint32_t aivId_{0};	71	uint32_t aivId_{0};
62	uint64_t aivNum_{0};	72	uint64_t aivNum_{0};
63	uint32_t round_{0};	73	uint32_t round_{0};
	@@ -95,10 +105,29 @@ private:
95	TBuf<> stateResetBuf_;	105	TBuf<> stateResetBuf_;
96	};	106	};
97		107
		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	+
98	template <TemplateTypeClass>	124	template <TemplateTypeClass>
99	__aicore__ inline void MTECommunication<TemplateType>::InitHcclContext()	125	__aicore__ inline void MTECommunication<TemplateType>::InitHcclContext()
100	{	126	{
101	hcclContext_ = (__gm__ Mc2Kernel::HcclOpParam*)GetHcclContext<HCCL_GROUP_ID_0>();	127	hcclContext_ = (__gm__ Mc2Kernel::HcclOpParam*)GetHcclContext<HCCL_GROUP_ID_0>();
		128	+ rankIdHccl_ = Mc2Kernel::GetRankId(hcclContext_);
		129	+ rankDimHccl_ = Mc2Kernel::GetRankDim(hcclContext_);
		130	+ isMc2Context_ = false;
102	}	131	}
103		132
104	template <TemplateTypeClass>	133	template <TemplateTypeClass>
	@@ -123,14 +152,14 @@ __aicore__ inline void MTECommunication<TemplateType>::InitGMTensor(GM_ADDR x, G
123	// =========== Win区相关 ===========	152	// =========== Win区相关 ===========
124	// Win区OOM检测适配，告知OOM框架Win区地址和大小	153	// Win区OOM检测适配，告知OOM框架Win区地址和大小
125	#if defined(ASCENDC_OOM) && ASCENDC_OOM == 1	154	#if defined(ASCENDC_OOM) && ASCENDC_OOM == 1
126	- for(uint64_t curRank = 0; curRank < hcclContext_->rankDim; ++curRank) {	155	+ for (uint64_t curRank = 0; curRank < rankDimHccl_; ++curRank) {
127	OOMCheckAddrRange(GetWinAddrGm(curRank), winSpaceGmSize);	156	OOMCheckAddrRange(GetWinAddrGm(curRank), winSpaceGmSize);
128	}	157	}
129	#endif	158	#endif
130		159
131	// 处理 0/1 分区标志位	160	// 处理 0/1 分区标志位
132	uint64_t currCoreFlagOffset = 2UL * SINGLE_STATE_REGION_SIZE + aivId_ * WIN_ADDR_ALIGN; // 计算当前核的标志位在Win区的偏移	161	uint64_t currCoreFlagOffset = 2UL * SINGLE_STATE_REGION_SIZE + aivId_ * WIN_ADDR_ALIGN; // 计算当前核的标志位在Win区的偏移
133	- selfWinFlagGMTensor_.SetGlobalBuffer((__gm__ uint32_t*)GetWinAddrGm(hcclContext_->rankId, currCoreFlagOffset));	162	+ selfWinFlagGMTensor_.SetGlobalBuffer((__gm__ uint32_t*)GetWinAddrGm(rankIdHccl_, currCoreFlagOffset));
134	LocalTensor<uint32_t> winFlagLocalTensor = winFlagsBuf_.Get<uint32_t>();	163	LocalTensor<uint32_t> winFlagLocalTensor = winFlagsBuf_.Get<uint32_t>();
135	DataCopy(winFlagLocalTensor, selfWinFlagGMTensor_, UB_ALIGN_BYTES / sizeof(uint32_t)); // GM -> UB	164	DataCopy(winFlagLocalTensor, selfWinFlagGMTensor_, UB_ALIGN_BYTES / sizeof(uint32_t)); // GM -> UB
136	SyncFunc<AscendC::HardEvent::MTE2_S>();	165	SyncFunc<AscendC::HardEvent::MTE2_S>();
	@@ -141,7 +170,7 @@ __aicore__ inline void MTECommunication<TemplateType>::InitGMTensor(GM_ADDR x, G
141		170
142	// 获取本卡地址写数据	171	// 获取本卡地址写数据
143	// 通过rankId和0/1分区标志位获取本地winIn区地址对应卡的数据区域	172	// 通过rankId和0/1分区标志位获取本地winIn区地址对应卡的数据区域
144	- GM_ADDR localDataSpaceGm = GetWinDataAddrGm(hcclContext_->rankId, winBufferFlags_);	173	+ GM_ADDR localDataSpaceGm = GetWinDataAddrGm(rankIdHccl_, winBufferFlags_);
145	localWinXGMTensor_.SetGlobalBuffer((__gm__ XType*)localDataSpaceGm);	174	localWinXGMTensor_.SetGlobalBuffer((__gm__ XType*)localDataSpaceGm);
146	localWinScaleGMTensor_.SetGlobalBuffer((__gm__ ScalesType*)(localDataSpaceGm + xSize)); // sclae数据跟在x后	175	localWinScaleGMTensor_.SetGlobalBuffer((__gm__ ScalesType*)(localDataSpaceGm + xSize)); // sclae数据跟在x后
147	}	176	}
	@@ -159,11 +188,12 @@ __aicore__ inline void MTECommunication<TemplateType>::InitBuffer(TPipe *tPipe)
159	tPipe->InitBuffer(xOutQueue_, BUFFER_NUM, xNumPerBlock_ * sizeof(OutputType)); // 用于输出的OutPutTensor	188	tPipe->InitBuffer(xOutQueue_, BUFFER_NUM, xNumPerBlock_ * sizeof(OutputType)); // 用于输出的OutPutTensor
160	tPipe->InitBuffer(winFlagsBuf_, UB_ALIGN_BYTES); // 用于读取0/1分区的标志位	189	tPipe->InitBuffer(winFlagsBuf_, UB_ALIGN_BYTES); // 用于读取0/1分区的标志位
161	tPipe->InitBuffer(writeStateBuf_, UB_ALIGN_BYTES); // 状态位每一个按32B对齐	190	tPipe->InitBuffer(writeStateBuf_, UB_ALIGN_BYTES); // 状态位每一个按32B对齐
162	- tPipe->InitBuffer(readStateBuf_, hcclContext_->rankDim * UB_ALIGN_BYTES); // 每次读 rankDim 个状态位	191	+ tPipe->InitBuffer(readStateBuf_, rankDimHccl_ * UB_ALIGN_BYTES); // 每次读 rankDim 个状态位
163	- tPipe->InitBuffer(stateResetBuf_, hcclContext_->rankDim * UB_ALIGN_BYTES); // 用于清理状态区	192	+ tPipe->InitBuffer(stateResetBuf_, rankDimHccl_ * UB_ALIGN_BYTES); // 用于清理状态区
164		193
165	stateResetTensor_ = stateResetBuf_.Get<float>();	194	stateResetTensor_ = stateResetBuf_.Get<float>();
166	- Duplicate<float>(stateResetTensor_, (float)0.0, static_cast<uint32_t>(hcclContext_->rankDim * FLOAT_UB_ALIGN_NUM)); // 用于状态区清零	195	+ Duplicate<float>(stateResetTensor_, (float)0.0,
		196	+ static_cast<uint32_t>(rankDimHccl_ * FLOAT_UB_ALIGN_NUM)); // 用于状态区清零
167	}	197	}
168		198
169	/**	199	/**
	@@ -269,7 +299,7 @@ __aicore__ inline void MTECommunication<TemplateType>::CopyDataToWin(uint64_t xS
269	}	299	}
270	if constexpr (isReduceScatter) {	300	if constexpr (isReduceScatter) {
271	// ReduceScatter过程，数据按卡均分，需要对卡进行遍历	301	// ReduceScatter过程，数据按卡均分，需要对卡进行遍历
272	- for(uint64_t curRank = 0; curRank < hcclContext_->rankDim; ++curRank) {	302	+ for (uint64_t curRank = 0; curRank < rankDimHccl_; ++curRank) {
273	// all2all过程，加上卡偏移	303	// all2all过程，加上卡偏移
274	uint64_t curRankXOffset = curXOffset + curRank * xSliceSizeNums;	304	uint64_t curRankXOffset = curXOffset + curRank * xSliceSizeNums;
275	uint64_t curRankScaleOffset = curScaleOffset + curRank * scaleSliceNums;	305	uint64_t curRankScaleOffset = curScaleOffset + curRank * scaleSliceNums;
	@@ -296,9 +326,9 @@ __aicore__ inline void MTECommunication<TemplateType>::CopyDataToWin(uint64_t xS
296	template <TemplateTypeClass>	326	template <TemplateTypeClass>
297	__aicore__ inline void MTECommunication<TemplateType>::WriteStatusToWin()	327	__aicore__ inline void MTECommunication<TemplateType>::WriteStatusToWin()
298	{	328	{
299	- uint32_t coreOffset = aivId_ * hcclContext_->rankDim; // Win区大小为 aivNum * rankDim, 此处计算核偏移	329	+ uint32_t coreOffset = aivId_ * rankDimHccl_; // Win区大小为 aivNum * rankDim, 此处计算核偏移
300	// 遍历每一张卡，给每一张卡都要写入状态	330	// 遍历每一张卡，给每一张卡都要写入状态
301	- for (uint32_t curRank = 0; curRank < hcclContext_->rankDim; ++curRank) {	331	+ for (uint32_t curRank = 0; curRank < rankDimHccl_; ++curRank) {
302	// 写入状态到对端，每个核写一个状态，表示自己的数据块已经写完	332	// 写入状态到对端，每个核写一个状态，表示自己的数据块已经写完
303	LocalTensor<float> statusTensor = writeStateBuf_.Get<float>();	333	LocalTensor<float> statusTensor = writeStateBuf_.Get<float>();
304	DataCopy<float>(statusTensor, stateResetTensor_, FLOAT_UB_ALIGN_NUM); // 先重置statusTensor数据，后面累加需要Tensor内全部数据，防止脏数据	334	DataCopy<float>(statusTensor, stateResetTensor_, FLOAT_UB_ALIGN_NUM); // 先重置statusTensor数据，后面累加需要Tensor内全部数据，防止脏数据
	@@ -308,7 +338,7 @@ __aicore__ inline void MTECommunication<TemplateType>::WriteStatusToWin()
308	GlobalTensor<float> stateGMTensor;	338	GlobalTensor<float> stateGMTensor;
309	stateGMTensor.SetGlobalBuffer((__gm__ float*)remoteWinStateGM);	339	stateGMTensor.SetGlobalBuffer((__gm__ float*)remoteWinStateGM);
310	// 不同卡上的核的状态写到相邻位置，读时可以一次读rankDim个状态, 状态区大小设计为 aivNum * ranDim	340	// 不同卡上的核的状态写到相邻位置，读时可以一次读rankDim个状态, 状态区大小设计为 aivNum * ranDim
311	- uint64_t curOffset = (coreOffset + hcclContext_->rankId) * FLOAT_UB_ALIGN_NUM; // 当前核偏移 + 卡偏移，按32B对齐	341	+ uint64_t curOffset = (coreOffset + rankIdHccl_) * FLOAT_UB_ALIGN_NUM; // 当前核偏移 + 卡偏移，按32B对齐
312	SyncFunc<AscendC::HardEvent::S_MTE3>();	342	SyncFunc<AscendC::HardEvent::S_MTE3>();
313	DataCopy(stateGMTensor[curOffset], statusTensor, FLOAT_UB_ALIGN_NUM); // 按32B对齐拷贝	343	DataCopy(stateGMTensor[curOffset], statusTensor, FLOAT_UB_ALIGN_NUM); // 按32B对齐拷贝
314	SyncFunc<AscendC::HardEvent::MTE3_S>();	344	SyncFunc<AscendC::HardEvent::MTE3_S>();
	@@ -325,16 +355,16 @@ __aicore__ inline void MTECommunication<TemplateType>::WriteStatusToWin()
325	template <TemplateTypeClass>	355	template <TemplateTypeClass>
326	__aicore__ inline void MTECommunication<TemplateType>::ReadStatus()	356	__aicore__ inline void MTECommunication<TemplateType>::ReadStatus()
327	{	357	{
328	- GM_ADDR stateGM = GetWinStatusAddrGm(hcclContext_->rankId, winBufferFlags_); // 获取本卡的状态区用于读取	358	+ GM_ADDR stateGM = GetWinStatusAddrGm(rankIdHccl_, winBufferFlags_); // 获取本卡的状态区用于读取
329	GlobalTensor<float> selfStatusWinTensor;	359	GlobalTensor<float> selfStatusWinTensor;
330	- uint32_t offset = aivId_ * hcclContext_->rankDim * FLOAT_UB_ALIGN_NUM; // 获取当前核所需读取状态位的头地址，状态按32B对齐	360	+ uint32_t offset = aivId_ * rankDimHccl_ * FLOAT_UB_ALIGN_NUM; // 获取当前核所需读取状态位的头地址，状态按32B对齐
331	selfStatusWinTensor.SetGlobalBuffer((__gm__ float*)(stateGM));	361	selfStatusWinTensor.SetGlobalBuffer((__gm__ float*)(stateGM));
332	LocalTensor<float> statusTensor = readStateBuf_.Get<float>();	362	LocalTensor<float> statusTensor = readStateBuf_.Get<float>();
333	float flag = 0; // 用于计算状态和	363	float flag = 0; // 用于计算状态和
334	- uint32_t statusCnt = hcclContext_->rankDim * FLOAT_UB_ALIGN_NUM; // 一次读rankDim个，按32B对齐	364	+ uint32_t statusCnt = rankDimHccl_ * FLOAT_UB_ALIGN_NUM; // 一次读rankDim个，按32B对齐
335	SumParams sumParams{1, statusCnt, statusCnt};	365	SumParams sumParams{1, statusCnt, statusCnt};
336	- float minTarget = hcclContext_->rankDim - (float)0.5;	366	+ float minTarget = rankDimHccl_ - (float)0.5;
337	- float maxTarget = hcclContext_->rankDim + (float)0.5;	367	+ float maxTarget = rankDimHccl_ + (float)0.5;
338	// 读取statusCnt个数据求和	368	// 读取statusCnt个数据求和
339	while ((flag < minTarget) \|\| (flag > maxTarget)) {	369	while ((flag < minTarget) \|\| (flag > maxTarget)) {
340	SyncFunc<AscendC::HardEvent::S_MTE2>();	370	SyncFunc<AscendC::HardEvent::S_MTE2>();
	@@ -375,6 +405,9 @@ __aicore__ inline void MTECommunication<TemplateType>::CopyResultToOutput(uint64
375	template <TemplateTypeClass>	405	template <TemplateTypeClass>
376	__aicore__ inline GM_ADDR MTECommunication<TemplateType>::GetWinAddrGm(uint32_t rankId, uint64_t offset)	406	__aicore__ inline GM_ADDR MTECommunication<TemplateType>::GetWinAddrGm(uint32_t rankId, uint64_t offset)
377	{	407	{
		408	+ if (isMc2Context_) {
		409	+ return (GM_ADDR)(mc2Context_->windowsIn[rankId] + offset);
		410	+ }
378	return (GM_ADDR)(hcclContext_->windowsIn[rankId] + offset);	411	return (GM_ADDR)(hcclContext_->windowsIn[rankId] + offset);
379	}	412	}
380		413
	@@ -387,6 +420,9 @@ __aicore__ inline GM_ADDR MTECommunication<TemplateType>::GetWinDataAddrGm(uint3
387	return GetWinAddrGm(rankId, STATE_WIN_SIZE);	420	return GetWinAddrGm(rankId, STATE_WIN_SIZE);
388	}	421	}
389	else {	422	else {
		423	+ if (isMc2Context_) {
		424	+ return (GM_ADDR)(mc2Context_->windowsOut[rankId]);
		425	+ }
390	// 若使用 1 分区，即WinOut	426	// 若使用 1 分区，即WinOut
391	return (GM_ADDR)(hcclContext_->windowsOut[rankId]);	427	return (GM_ADDR)(hcclContext_->windowsOut[rankId]);
392	}	428	}
	@@ -405,5 +441,7 @@ __aicore__ inline GM_ADDR MTECommunication<TemplateType>::GetWinStatusAddrGm(uin
405	return GetWinAddrGm(rankId, SINGLE_STATE_REGION_SIZE);	441	return GetWinAddrGm(rankId, SINGLE_STATE_REGION_SIZE);
406	}	442	}
407	}	443	}
		444	+
408	} // QuantMTECommImpl	445	} // QuantMTECommImpl
		446	+
409	#endif // MTE_COMMON_H	447	#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
30	template<uint32_t quantReduceScatterCommMode>	30	template<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(x, scales, output, &pipe, &tilingData);	40	+ op.Init(nullptr, 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,6 +35,7 @@ namespace QuantReduceScatterImpl {
35	using namespace QuantMTECommImpl;	35	using namespace QuantMTECommImpl;
36	using namespace VectorComputeImpl;	36	using namespace VectorComputeImpl;
37	using namespace AscendC;	37	using namespace AscendC;
		38	+using namespace Mc2Aclnn;
38		39
39	constexpr static uint64_t MX_SCALES_LAST_DIM = 2U; // MX量化scales最后一维的大小	40	constexpr static uint64_t MX_SCALES_LAST_DIM = 2U; // MX量化scales最后一维的大小
40		41
	@@ -42,7 +43,7 @@ template<TemplateTypeClass>
42	class QuantReduceScatterMte {	43	class QuantReduceScatterMte {
43	public:	44	public:
44	__aicore__ inline QuantReduceScatterMte() {};	45	__aicore__ inline QuantReduceScatterMte() {};
45	- __aicore__ inline void Init(GM_ADDR x, GM_ADDR scales, GM_ADDR output,	46	+ __aicore__ inline void Init(GM_ADDR mc2Context, GM_ADDR x, GM_ADDR scales, GM_ADDR output,
46	TPipe pipe, const QuantReduceScatterTilingData tilingData);	47	TPipe pipe, const QuantReduceScatterTilingData tilingData);
47	__aicore__ inline void Process();	48	__aicore__ inline void Process();
48	private:	49	private:
	@@ -86,10 +87,10 @@ private:
86	};	87	};
87		88
88	template <TemplateTypeClass>	89	template <TemplateTypeClass>
89	-__aicore__ inline void QuantReduceScatterMte<TemplateType>::Init(GM_ADDR x, GM_ADDR scales,	90	+__aicore__ inline void QuantReduceScatterMte<TemplateType>::Init(GM_ADDR mc2Context, GM_ADDR x, GM_ADDR scales,
90	GM_ADDR output, TPipe tPipe, const QuantReduceScatterTilingData tilingData)	91	GM_ADDR output, TPipe tPipe, const QuantReduceScatterTilingData tilingData)
91	{	92	{
92	- mteComm_.InitHcclContext();	93	+ mteComm_.InitMc2Context(mc2Context, tilingData);
93	ParseTilingInfo(tilingData);	94	ParseTilingInfo(tilingData);
94	tPipe->Reset();	95	tPipe->Reset();
95	ComputeXPerBlock(tilingData, tPipe);	96	ComputeXPerBlock(tilingData, tPipe);
	@@ -109,9 +110,9 @@ __aicore__ inline void QuantReduceScatterMte<TemplateType>::ParseTilingInfo(
109	if constexpr(AscendC::IsSameType<ScalesType, fp8_e8m0_t>::value) {	110	if constexpr(AscendC::IsSameType<ScalesType, fp8_e8m0_t>::value) {
110	scaleSize_ *= MX_SCALES_LAST_DIM;	111	scaleSize_ *= MX_SCALES_LAST_DIM;
111	}	112	}
112	- uint64_t xSliceSize = xSize_ / (mteComm_.hcclContext_->rankDim);	113	+ uint64_t xSliceSize = xSize_ / (mteComm_.rankDimHccl_);
113	xSliceSizeNums_ = xSliceSize / sizeof(XType);	114	xSliceSizeNums_ = xSliceSize / sizeof(XType);
114	- scaleSliceNums_ = scaleSize_ / (mteComm_.hcclContext_->rankDim * sizeof(ScalesType));	115	+ scaleSliceNums_ = scaleSize_ / (mteComm_.rankDimHccl_ * sizeof(ScalesType));
115	}	116	}
116		117
117	template <TemplateTypeClass>	118	template <TemplateTypeClass>
	@@ -124,8 +125,8 @@ __aicore__ inline void QuantReduceScatterMte<TemplateType>::ComputeXPerBlock(
124	uint64_t mteCommFixedSpace = BUFFER_NUM * X_BLOCK_BYTES + // scaleQueue_	125	uint64_t mteCommFixedSpace = BUFFER_NUM * X_BLOCK_BYTES + // scaleQueue_
125	UB_ALIGN_BYTES + // winFlagsBuf_	126	UB_ALIGN_BYTES + // winFlagsBuf_
126	UB_ALIGN_BYTES + // writeStateBuf_	127	UB_ALIGN_BYTES + // writeStateBuf_
127	- mteComm_.hcclContext_->rankDim * UB_ALIGN_BYTES + // readStateBuf_	128	+ mteComm_.rankDimHccl_ * UB_ALIGN_BYTES + // readStateBuf_
128	- mteComm_.hcclContext_->rankDim * UB_ALIGN_BYTES; // stateResetBuf_	129	+ mteComm_.rankDimHccl_ * UB_ALIGN_BYTES; // stateResetBuf_
129		130
130	// 动态开销：每增加 1 个 x 需要的 UB 字节（整数部分，分数部分见下方比例校正）	131	// 动态开销：每增加 1 个 x 需要的 UB 字节（整数部分，分数部分见下方比例校正）
131	uint64_t baseDynamic = BUFFER_NUM * sizeof(OutputType) + // xOutQueue_	132	uint64_t baseDynamic = BUFFER_NUM * sizeof(OutputType) + // xOutQueue_
	@@ -239,7 +240,6 @@ __aicore__ inline void QuantReduceScatterMte<TemplateType>::ReadDataBlockReduceS
239	scaleInQue_.FreeTensor(scaleTmpTensor);	240	scaleInQue_.FreeTensor(scaleTmpTensor);
240	}	241	}
241		242
242	-
243	template <TemplateTypeClass>	243	template <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_.hcclContext_->rankId;	278	+ uint32_t startRankId = mteComm_.rankIdHccl_;
279	- for (uint32_t i = 0; i < mteComm_.hcclContext_->rankDim; ++i) {	279	+ for (uint32_t i = 0; i < mteComm_.rankDimHccl_; ++i) {
280	- uint32_t remoteRankId = (startRankId + i) % mteComm_.hcclContext_->rankDim;	280	+ uint32_t remoteRankId = (startRankId + i) % mteComm_.rankDimHccl_;
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_.hcclContext_->rankId * xSliceSizeNums_;	290	+ uint64_t curRankXOffset = curXOffset + mteComm_.rankIdHccl_ * xSliceSizeNums_;
291	- uint64_t curRankScaleOffset = curScaleOffset + mteComm_.hcclContext_->rankId * scaleSliceNums_;	291	+ uint64_t curRankScaleOffset = curScaleOffset + mteComm_.rankIdHccl_ * 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+1-0

	@@ -27,6 +27,7 @@ struct QuantReduceScatterTilingInfo {
27	uint64_t totalWinSize; // Win区总大小，即HCCL_BUFFER_SIZE	27	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;
30	};	31	};
31		32
32	struct QuantReduceScatterTilingData {	33	struct QuantReduceScatterTilingData {

Mmc2/quant_reduce_scatter/tests/ut/op_host/arch35/test_quant_reduce_scatter_tiling.cpp+14-14

Amc2/quant_reduce_scatter_v2/CMakeLists.txt+21-0

@@ -0,0 +1,21 @@
	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)

Amc2/quant_reduce_scatter_v2/op_api/.gitkeep+0-0

The file is empty

Amc2/quant_reduce_scatter_v2/op_graph/CMakeLists.txt+14-0

@@ -0,0 +1,14 @@
	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()

Amc2/quant_reduce_scatter_v2/op_graph/fusion_pass/.gitkeep+0-0

The file is empty

Amc2/quant_reduce_scatter_v2/op_graph/quant_reduce_scatter_v2_proto.h+58-0

@@ -0,0 +1,58 @@
	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_

Amc2/quant_reduce_scatter_v2/op_host/CMakeLists.txt+35-0

@@ -0,0 +1,35 @@
	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()

Amc2/quant_reduce_scatter_v2/op_host/op_tiling/quant_reduce_scatter_v2_tiling.cpp+171-0

@@ -0,0 +1,171 @@
	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

Amc2/quant_reduce_scatter_v2/op_host/quant_reduce_scatter_v2_def.cpp+94-0

@@ -0,0 +1,94 @@
	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

Amc2/quant_reduce_scatter_v2/op_kernel/quant_reduce_scatter_v2.cpp+51-0

@@ -0,0 +1,51 @@
	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	+}

Amc2/quant_reduce_scatter_v2/tests/CMakeLists.txt+16-0

@@ -0,0 +1,16 @@
	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()

Amc2/quant_reduce_scatter_v2/tests/ut/CMakeLists.txt+16-0

@@ -0,0 +1,16 @@
	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()

Amc2/quant_reduce_scatter_v2/tests/ut/op_api/CMakeLists.txt+13-0

@@ -0,0 +1,13 @@
	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()

Amc2/quant_reduce_scatter_v2/tests/ut/op_api/test_aclnn_quant_reduce_scatter_v2.cpp+9-0

@@ -0,0 +1,9 @@
	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	+**/

Amc2/quant_reduce_scatter_v2/tests/ut/op_host/CMakeLists.txt+22-0

@@ -0,0 +1,22 @@
	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()

Amc2/quant_reduce_scatter_v2/tests/ut/op_host/arch35/CMakeLists.txt+13-0

@@ -0,0 +1,13 @@
	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()

Amc2/quant_reduce_scatter_v2/tests/ut/op_host/arch35/test_quant_reduce_scatter_v2_tiling.cpp+9-0

@@ -0,0 +1,9 @@
	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+19-0

	@@ -1429,6 +1429,25 @@ mc2:
1429	examples: False	1429	examples: False
1430	options:	1430	options:
1431	- quant_reduce_scatter	1431	- 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
1432		1451
1433	moe_distribute:	1452	moe_distribute:
1434	distribute_barrier:	1453	distribute_barrier:

	@@ -67,7 +67,8 @@ static ge::graphStatus SetHcommCfg(const gert::TilingContext *context, QuantAllR
67	* @param tilingData: 框架根据context的opName匹配tiling模板，计算产生的tilingData	67	* @param tilingData: 框架根据context的opName匹配tiling模板，计算产生的tilingData
68	* @return	68	* @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)
71	{	72	{
72	fe::PlatFormInfos *platformInfoPtr = context->GetPlatformInfo();	73	fe::PlatFormInfos *platformInfoPtr = context->GetPlatformInfo();
73	platform_ascendc::PlatformAscendC ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);	74	platform_ascendc::PlatformAscendC ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
	@@ -75,22 +76,25 @@ static void SetTilingData(gert::TilingContext *context, QuantAllReduceTilingData
75	uint32_t aivNum = ascendcPlatform.GetCoreNumAiv();	76	uint32_t aivNum = ascendcPlatform.GetCoreNumAiv();
76	context->SetBlockDim(ascendcPlatform.CalcTschBlockDim(aivNum, 0, aivNum));	77	context->SetBlockDim(ascendcPlatform.CalcTschBlockDim(aivNum, 0, aivNum));
77	tilingData.quantAllReduceTilingInfo.aivNum = aivNum;	78	tilingData.quantAllReduceTilingInfo.aivNum = aivNum;
78	- uint64_t xValueBS = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);	79	+ uint64_t xValueBS = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ZERO);
79	- uint64_t xValueH = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ONE);	80	+ uint64_t xValueH = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
80	- uint64_t scalesValueH = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(DIM_ONE);	81	+ uint64_t scalesValueH = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_ONE);
81	// context->GetInputShape在函数CheckInputTensorDim中已经校验	82	// context->GetInputShape在函数CheckInputTensorDim中已经校验
82	- if (context->GetInputShape(X_INDEX)->GetStorageShape().GetDimNum() == THREE_DIMS) {	83	+ if (context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDimNum() == THREE_DIMS) {
83	- xValueBS = xValueBS * context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_ONE);	84	+ xValueBS = xValueBS * context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_ONE);
84	- xValueH = context->GetInputShape(X_INDEX)->GetStorageShape().GetDim(DIM_TWO);	85	+ xValueH = context->GetInputShape(config.X_INDEX)->GetStorageShape().GetDim(DIM_TWO);
85	- scalesValueH = context->GetInputShape(SCALES_INDEX)->GetStorageShape().GetDim(DIM_TWO);	86	+ scalesValueH = context->GetInputShape(config.SCALES_INDEX)->GetStorageShape().GetDim(DIM_TWO);
86	}	87	}
87	tilingData.quantAllReduceTilingInfo.bs = xValueBS;	88	tilingData.quantAllReduceTilingInfo.bs = xValueBS;
88	tilingData.quantAllReduceTilingInfo.hiddenSize = xValueH;	89	tilingData.quantAllReduceTilingInfo.hiddenSize = xValueH;
89	tilingData.quantAllReduceTilingInfo.scaleHiddenSize = scalesValueH;	90	tilingData.quantAllReduceTilingInfo.scaleHiddenSize = scalesValueH;
90	tilingData.quantAllReduceTilingInfo.totalWinSize = mc2tiling::Mc2TilingUtils::GetMaxWindowSize();	91	tilingData.quantAllReduceTilingInfo.totalWinSize = mc2tiling::Mc2TilingUtils::GetMaxWindowSize();
		92	+ tilingData.quantAllReduceTilingInfo.isMc2Context = config.isMc2Context;
91	}	93	}
92		94
93	-// 基于 TARGET_ITER 公式计算 host 推荐的 xPerBlock，写入 tilingData	95	+/**
		96	+ * @brief 基于 TARGET_ITER 公式计算 host 推荐的 xPerBlock，写入 tilingData
		97	+ */
94	static void SetXPerBlock(QuantAllReduceTilingData &tilingData)	98	static void SetXPerBlock(QuantAllReduceTilingData &tilingData)
95	{	99	{
96	constexpr uint32_t TARGET_ITER = 3U; // T=3 命中 DoubleBuffer 甜点	100	constexpr uint32_t TARGET_ITER = 3U; // T=3 命中 DoubleBuffer 甜点
	@@ -141,12 +145,16 @@ static ge::graphStatus QuantAllReduceTilingFunc(gert::TilingContext *context)
141	OP_TILING_CHECK(QuantReduceScatterUtilTiling::CheckNpuArch(context) != ge::GRAPH_SUCCESS,	145	OP_TILING_CHECK(QuantReduceScatterUtilTiling::CheckNpuArch(context) != ge::GRAPH_SUCCESS,
142	OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(nodeName, "npuArch", "non-DAV_3510", "The value of npuArch must be DAV_3510"),	146	OP_LOGE_FOR_INVALID_VALUE_WITH_REASON(nodeName, "npuArch", "non-DAV_3510", "The value of npuArch must be DAV_3510"),
143	return ge::GRAPH_FAILED);	147	return ge::GRAPH_FAILED);
144	- OP_TILING_CHECK(QuantReduceScatterUtilTiling::CheckTilingFunc(context, runInfo, OpType::OP_QUANT_ALL_REDUCE) !=	148	+ QuantReduceScatterConfig config;
145	- ge::GRAPH_SUCCESS,	149	+ config.X_INDEX = 0;
		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,
146	OP_LOGE(nodeName, "tiling check failed in quant_all_reduce."), return ge::GRAPH_FAILED);	154	OP_LOGE(nodeName, "tiling check failed in quant_all_reduce."), return ge::GRAPH_FAILED);
147	OP_TILING_CHECK(SetHcommCfg(context, tilingData, runInfo) != ge::GRAPH_SUCCESS,	155	OP_TILING_CHECK(SetHcommCfg(context, tilingData, runInfo) != ge::GRAPH_SUCCESS,
148	OP_LOGE(nodeName, "SetHCommCfg failed."), return ge::GRAPH_FAILED);	156	OP_LOGE(nodeName, "SetHCommCfg failed."), return ge::GRAPH_FAILED);
149	- SetTilingData(context, *tilingData);	157	+ SetTilingData(context, *tilingData, config);
150	SetXPerBlock(*tilingData);	158	SetXPerBlock(*tilingData);
151	SetTilingKey(context);	159	SetTilingKey(context);
152	PrintTilingDataInfo(context, *tilingData);	160	PrintTilingDataInfo(context, *tilingData);

	@@ -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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	68	+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 80 64 314572800 4398046514176 0 ", {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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	74	+ ge::GRAPH_SUCCESS, 0UL, "2048 5120 40 64 314572800 4398046517248 0 ", {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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	80	+ ge::GRAPH_SUCCESS, 0UL, "1024 7168 56 64 314572800 4398046515200 0 ", {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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	86	+ ge::GRAPH_SUCCESS, 0UL, "1024 4096 64 64 314572800 4398046513152 0 ", {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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	92	+ ge::GRAPH_SUCCESS, 0UL, "2048 4096 32 64 314572800 4398046516224 0 ", {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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	98	+ ge::GRAPH_SUCCESS, 0UL, "1024 8192 64 64 314572800 4398046516224 0 ", {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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	267	+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 80 64 314572800 4398046524416 0 ", {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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	273	+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 40 64 314572800 4398046517248 0 ", {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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	279	+ ge::GRAPH_SUCCESS, 0UL, "1024 1024 8 64 314572800 4398046513152 0 ", {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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	285	+ ge::GRAPH_SUCCESS, 0UL, "1024 8192 128 64 314572800 4398046516224 0 ", {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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	291	+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 80 64 314572800 4398046514176 0 ", {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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	297	+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 40 64 314572800 4398046514176 0 ", {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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	303	+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 40 64 314572800 4398046514176 0 ", {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 ", {16777216}, MC2_TILING_DATA_RESERVED_LEN},	309	+ ge::GRAPH_SUCCESS, 0UL, "1024 5120 40 64 314572800 4398046524416 0 ", {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,