mhc_post_backward A2/A3支持-ops-transformer-AtomGit

已合并

mhc_post_backward A2/A3支持 #4377

fazhenyao123创建于 4月20日从fazhenyao123/ops-transformer： 4_2 合入到 cann/ops-transformer： master

mhc_post_backward A2/A3支持 #4377

已合并

fazhenyao123创建于 4月20日

从fazhenyao123/ops-transformer 4_2合入到cann/ops-transformer master

共 7 个文件变更+843-0

Amhc/mhc_post_backward/CMakeLists.txt+19-0

@@ -0,0 +1,19 @@
	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	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	12	+if(NOT ENABLE_TEST AND NOT BENCHMARK)
	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()

Amhc/mhc_post_backward/op_host/CMakeLists.txt+26-0

@@ -0,0 +1,26 @@
	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 (BUILD_OPEN_PROJECT)
	12	+ target_sources(op_host_aclnn PRIVATE
	13	+ mhc_post_backward_def.cpp
	14	+ )
	15	+ add_ops_compile_options(
	16	+ OP_NAME MhcPostBackward
	17	+ OPTIONS --cce-auto-sync=off
	18	+ -Wno-deprecated-declarations
	19	+ -Werror
	20	+ )
	21	+endif()
	22	+
	23	+if(NOT BUILD_OPS_RTY_KERNEL)
	24	+ add_op_to_compiled_list()
	25	+ add_modules_sources(OPTYPE mhc_post_backward ACLNNTYPE aclnn)
	张张张硕累4月20日【review】不建议使用自动生成ACLNN，建议新增手写aclnn接口
	26	+endif()

Amhc/mhc_post_backward/op_host/mhc_post_backward_def.cpp+93-0

@@ -0,0 +1,93 @@
	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 mhc_post_backward.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+ #include "register/op_def_registry.h"
	17	+
	18	+namespace ops{
	19	+class MhcPostBackward : public OpDef
	20	+{
	21	+public:
	22	+ explicit MhcPostBackward(const char* name) : OpDef(name)
	23	+ {
	24	+ this->Input("grad_y")
	25	+ .ParamType(REQUIRED)
	26	+ .DataType({ge::DT_BF16, ge::DT_FLOAT16})
	27	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	28	+ .UnknownShapeFormat(
	29	+ {ge::FORMAT_ND, ge::FORMAT_ND})
	30	+ .AutoContiguous();
	31	+ this->Input("x")
	32	+ .ParamType(REQUIRED)
	33	+ .DataType({ge::DT_BF16, ge::DT_FLOAT16})
	34	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	35	+ .UnknownShapeFormat(
	36	+ {ge::FORMAT_ND, ge::FORMAT_ND})
	37	+ .AutoContiguous();
	38	+ this->Input("h_res")
	39	+ .ParamType(REQUIRED)
	40	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT})
	41	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	42	+ .UnknownShapeFormat(
	43	+ {ge::FORMAT_ND, ge::FORMAT_ND})
	44	+ .AutoContiguous();
	45	+ this->Input("h_out")
	46	+ .ParamType(REQUIRED)
	47	+ .DataType({ge::DT_BF16, ge::DT_FLOAT16})
	48	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	49	+ .UnknownShapeFormat(
	50	+ {ge::FORMAT_ND, ge::FORMAT_ND})
	51	+ .AutoContiguous();
	52	+ this->Input("h_post")
	53	+ .ParamType(REQUIRED)
	54	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT})
	55	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	56	+ .UnknownShapeFormat(
	57	+ {ge::FORMAT_ND, ge::FORMAT_ND})
	58	+ .AutoContiguous();
	59	+ this->Output("grad_x")
	60	+ .ParamType(REQUIRED)
	61	+ .DataType({ge::DT_BF16, ge::DT_FLOAT16})
	62	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	63	+ .UnknownShapeFormat(
	64	+ {ge::FORMAT_ND, ge::FORMAT_ND})
	65	+ .AutoContiguous();
	66	+ this->Output("grad_h_res")
	67	+ .ParamType(REQUIRED)
	68	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT})
	69	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	70	+ .UnknownShapeFormat(
	71	+ {ge::FORMAT_ND, ge::FORMAT_ND})
	72	+ .AutoContiguous();
	73	+ this->Output("grad_h_out")
	74	+ .ParamType(REQUIRED)
	75	+ .DataType({ge::DT_BF16, ge::DT_FLOAT16})
	76	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	77	+ .UnknownShapeFormat(
	78	+ {ge::FORMAT_ND, ge::FORMAT_ND})
	79	+ .AutoContiguous();
	80	+ this->Output("grad_h_post")
	81	+ .ParamType(REQUIRED)
	82	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT})
	83	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND})
	84	+ .UnknownShapeFormat(
	85	+ {ge::FORMAT_ND, ge::FORMAT_ND})
	86	+ .AutoContiguous();
	87	+ this->AICore().AddConfig("ascend910b");
	88	+ this->AICore().AddConfig("ascend910_93");
	张张张硕累4月20日此条代码评论区间+87至+88 【review】此处无ascend950，与cmake不符，需要确认是否支持950芯片，若支持，此处应当新增this->AICore().AddConfig("ascend950"); fazhenyao123 4月20日评论： 4月20日评论：
	89	+ }
	90	+};
	91	+
	92	+OP_ADD(MhcPostBackward);
	93	+}//namespace ops
	CANN-robot4月20日代码结构与可维护性: 文件末尾缺少换行符（No newline at end of file）。虽然这不影响编译，但不符合许多代码规范（如POSIX标准），可能导致某些文本处理工具（如diff、cat、某些IDE）显示警告或行为不一致。在版本控制系统中，缺少结尾换行符可能引起不必要的diff噪音。问题类型: 代码结构与可维护性文件路径: `mhc/mhc_post_backward/op_host/mhc_post_backward_def.cpp` 行号: 93 问题代码: }//namespace ops 修改建议: 在文件最后一行添加一个换行符（空行）。 --- 此评论由代码审查工具自动生成

Amhc/mhc_post_backward/op_host/mhc_post_backward_tiling.cpp+185-0

@@ -0,0 +1,185 @@
	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 mhc_post_backward_tiling.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+#include "mhc_post_backward_tiling.h"
	17	+#include "log/log.h"
	18	+#include "register/op_def_registry.h"
	19	+#include "tiling/tiling_api.h"
	20	+#include "tiling/platform/platform_ascendc.h"
	21	+
	22	+using namespace ge;
	23	+using namespace std;
	24	+using namespace AscendC;
	25	+
	26	+namespace {
	27	+constexpr uint8_t GRAD_Y_IDX = 0;
	28	+constexpr uint8_t X_IDX = 1;
	29	+constexpr uint8_t H_RES_IDX = 2;
	30	+constexpr uint8_t H_OUT_IDX = 3;
	31	+constexpr uint8_t H_POST_IDX = 4;
	32	+
	33	+constexpr uint8_t X_MIX_GRAD_IDX = 0;
	34	+constexpr uint8_t H_MIX_GRAD_IDX = 1;
	35	+
	36	+constexpr uint8_t SIZE_BFLOAT16 = 2;
	37	+constexpr uint8_t SIZE_FLOAT = 4;
	38	+
	39	+static int32_t GetCeilInt(int32_t value1, int32_t value2)
	40	+{
	41	+ if (value2 == 0) {
	42	+ return value1;
	43	+ }
	44	+ return static_cast<int32_t>((value1 + value2 - 1) / value2);
	45	+}
	46	+
	47	+}
	48	+
	49	+namespace optiling {
	50	+namespace mhc_post_backward {
	51	+
	52	+const uint32_t BLOCK_C = 1024;
	53	+
	54	+ge::graphStatus TilingCompute(gert::TilingContext *context)
	55	+{
	56	+ if (context == nullptr) {
	57	+ return ge::GRAPH_FAILED;
	58	+ }
	59	+ MhcPostBackwardTilingData tiling;
	60	+ auto platformInfoptr = context->GetPlatformInfo();
	61	+ if (platformInfoptr == nullptr) {
	62	+ return ge::GRAPH_FAILED;
	63	+ }
	64	+
	65	+ auto ascendplatformInfo = platform_ascendc::PlatformAscendC(platformInfoptr);
	66	+ const auto coreNumber = ascendplatformInfo.GetCoreNumAiv();
	67	+
	68	+ auto gradYTensor = context->GetInputTensor(GRAD_Y_IDX);
	69	+ OP_CHECK_NULL_WITH_CONTEXT(context, gradYTensor);
	70	+ auto xTensor = context->GetInputTensor(X_IDX);
	71	+ OP_CHECK_NULL_WITH_CONTEXT(context, xTensor);
	72	+ auto hResTensor = context->GetInputTensor(H_RES_IDX);
	73	+ OP_CHECK_NULL_WITH_CONTEXT(context, hResTensor);
	74	+ auto hOutTensor = context->GetInputTensor(H_OUT_IDX);
	75	+ OP_CHECK_NULL_WITH_CONTEXT(context, hOutTensor);
	76	+ auto hPostTensor = context->GetInputTensor(H_POST_IDX);
	77	+ OP_CHECK_NULL_WITH_CONTEXT(context, hPostTensor);
	78	+
	79	+ auto gradYDesc = context->GetInputDesc(GRAD_Y_IDX);
	80	+ OP_CHECK_NULL_WITH_CONTEXT(context, gradYDesc);
	81	+ auto gradYDtype = gradYDesc->GetDataType();
	82	+ OP_CHECK_IF(
	83	+ gradYDtype != ge::DataType::DT_BF16 && gradYDtype != ge::DataType::DT_FLOAT16,
	84	+ OP_LOGE(context->GetNodeName(), "grad_y dtype only supports bf16,half."),
	85	+ return ge::GRAPH_FAILED);
	86	+
	87	+ auto xDesc = context->GetInputDesc(X_IDX);
	88	+ OP_CHECK_NULL_WITH_CONTEXT(context, xDesc);
	89	+ OP_CHECK_IF(
	90	+ xDesc->GetDataType() != gradYDtype,
	91	+ OP_LOGE(context->GetNodeName(), "the dtype of x should be same with grad_y."),
	92	+ return ge::GRAPH_FAILED);
	93	+
	94	+ auto hOutDesc = context->GetInputDesc(H_OUT_IDX);
	95	+ OP_CHECK_NULL_WITH_CONTEXT(context, hOutDesc);
	96	+ OP_CHECK_IF(
	97	+ hOutDesc->GetDataType() != gradYDtype,
	98	+ OP_LOGE(context->GetNodeName(), "the dtype of h_out should be same with grad_y."),
	99	+ return ge::GRAPH_FAILED);
	100	+
	101	+ auto hResDesc = context->GetInputDesc(H_RES_IDX);
	102	+ OP_CHECK_NULL_WITH_CONTEXT(context, hResDesc);
	103	+ OP_CHECK_IF(
	104	+ hResDesc->GetDataType() != ge::DataType::DT_FLOAT,
	105	+ OP_LOGE(context->GetNodeName(), "h_res dtype only supports float32."),
	106	+ return ge::GRAPH_FAILED);
	107	+
	108	+ auto hPostDesc = context->GetInputDesc(H_POST_IDX);
	109	+ OP_CHECK_NULL_WITH_CONTEXT(context, hPostDesc);
	110	+ OP_CHECK_IF(
	111	+ hPostDesc->GetDataType() != ge::DataType::DT_FLOAT,
	112	+ OP_LOGE(context->GetNodeName(), "h_post dtype only supports float32."),
	113	+ return ge::GRAPH_FAILED);
	114	+
	115	+ const auto dFPostResShape = gradYTensor->GetStorageShape();
	ww4月23日未校验维度数量和值合法性 fazhenyao123 4月25日评论： 4月25日评论：
	116	+
	117	+ const uint32_t totalTasks = dFPostResShape.GetDim(0);
	118	+
	119	+ uint64_t frontCore = totalTasks % coreNumber != 0 ? static_cast<uint64_t>(totalTasks % coreNumber) : coreNumber;
	120	+ uint64_t tailCore = totalTasks <= coreNumber ? 0 : coreNumber - frontCore;
	121	+
	122	+ int32_t singleCoreBS = GetCeilInt(totalTasks, coreNumber);
	123	+ int32_t tailBS = totalTasks / coreNumber;
	124	+
	125	+ const int32_t coreUsed = frontCore + tailCore;
	126	+ uint64_t ubSizePlatForm;
	127	+ ascendplatformInfo.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatForm);
	128	+
	129	+ uint32_t dFPostResSize = gradYTensor->GetShapeSize();
	130	+ uint32_t xSize = xTensor->GetShapeSize();
	131	+ uint32_t hResSize = hResTensor->GetShapeSize();
	132	+ uint32_t hOutSize = hOutTensor->GetShapeSize();
	133	+ uint32_t hPostSize = hPostTensor->GetShapeSize();
	134	+
	135	+ const uint32_t n = dFPostResShape.GetDim(1);
	136	+ const uint32_t alignN = GetCeilInt(n * SIZE_FLOAT, 32) * 32 /SIZE_FLOAT;
	ww4月23日 n 为 uint32_t，SIZE_FLOAT 为 4，乘法在 uint32_t 空间计算若 n 接近 UINT32_MAX/4，乘法可能溢出回绕 fazhenyao123 4月23日评论： 4月23日评论：
	137	+ const uint32_t channel = dFPostResShape.GetDim(2);
	138	+ const uint32_t blockChannel = BLOCK_C > channel ? channel : BLOCK_C;
	139	+ const uint32_t loopC = channel / blockChannel;
	140	+ const uint32_t tailC = channel % blockChannel;
	141	+
	142	+ context->SetBlockDim(coreUsed);
	143	+
	144	+ tiling.set_singleCoreBS(singleCoreBS);
	145	+ tiling.set_tailBS(tailBS);
	146	+ tiling.set_coreUsed(coreUsed);
	147	+ tiling.set_frontCore(frontCore);
	148	+ tiling.set_tailCore(tailCore);
	149	+
	150	+ tiling.set_dFPostResSize(dFPostResSize);
	151	+ tiling.set_xSize(xSize);
	152	+ tiling.set_hResSize(hResSize);
	153	+ tiling.set_hOutSize(hOutSize);
	154	+ tiling.set_hPostSize(hPostSize);
	155	+
	156	+ tiling.set_channel(channel);
	157	+ tiling.set_blockChannel(blockChannel);
	158	+ tiling.set_n(n);
	159	+ tiling.set_alignN(alignN);
	160	+ tiling.set_tailC(tailC);
	161	+ tiling.set_loopC(loopC);
	162	+
	163	+ tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
	164	+ context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());
	165	+
	张张张硕累4月21日此条代码评论区间+163至+165 【review】tiling里没设置workspace大小，会导致workspace大小为随机值，可能导致模型里显存问题，需要设置sysworkspace
	166	+ size_t* workspaces = context->GetWorkspaceSizes(1);
	167	+ workspaces[0] = ascendplatformInfo.GetLibApiWorkSpaceSize();
	168	+ return ge::GRAPH_SUCCESS;
	169	+}
	170	+
	171	+ge::graphStatus Tiling4MhcPostBackward(gert::TilingContext *context)
	172	+{
	173	+ return TilingCompute(context);
	174	+}
	175	+ge::graphStatus TilingPrepareForMhcPostBackward(gert::TilingParseContext* context)
	176	+{
	177	+ (void)context;
	178	+ return ge::GRAPH_SUCCESS;
	179	+}
	180	+
	181	+IMPL_OP_OPTILING(MhcPostBackward)
	182	+ .Tiling(Tiling4MhcPostBackward)
	183	+ .TilingParse<MhcPostBackwardCompileInfo>(TilingPrepareForMhcPostBackward);
	184	+}
	185	+}

Amhc/mhc_post_backward/op_host/mhc_post_backward_tiling.h+50-0

@@ -0,0 +1,50 @@
	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 mhc_post_backward_tiling.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#ifndef MHC_POST_BACKWARD_TILING_H
	17	+#define MHC_POST_BACKWARD_TILING_H
	18	+
	19	+#include "register/tilingdata_base.h"
	20	+
	21	+namespace optiling {
	22	+BEGIN_TILING_DATA_DEF(MhcPostBackwardTilingData)
	L L liuguoyue4月23日需要使用新的tiling注册方式 fazhenyao123 4月25日评论： 4月25日评论：
	23	+ TILING_DATA_FIELD_DEF(uint64_t, coreUsed);
	24	+ TILING_DATA_FIELD_DEF(uint64_t, frontCore);
	25	+ TILING_DATA_FIELD_DEF(uint64_t, tailCore);
	26	+ TILING_DATA_FIELD_DEF(uint64_t, singleCoreBS);
	27	+ TILING_DATA_FIELD_DEF(uint64_t, tailBS);
	28	+ TILING_DATA_FIELD_DEF(uint64_t, dFPostResSize);
	29	+ TILING_DATA_FIELD_DEF(uint64_t, xSize);
	30	+ TILING_DATA_FIELD_DEF(uint64_t, hResSize);
	31	+ TILING_DATA_FIELD_DEF(uint64_t, hOutSize);
	32	+ TILING_DATA_FIELD_DEF(uint64_t, hPostSize);
	33	+ TILING_DATA_FIELD_DEF(uint64_t, channel);
	34	+ TILING_DATA_FIELD_DEF(uint64_t, blockChannel);
	35	+ TILING_DATA_FIELD_DEF(uint64_t, n);
	36	+ TILING_DATA_FIELD_DEF(uint64_t, alignN);
	37	+ TILING_DATA_FIELD_DEF(uint64_t, tileC);
	38	+ TILING_DATA_FIELD_DEF(uint64_t, tailC);
	39	+ TILING_DATA_FIELD_DEF(uint64_t, loopC);
	40	+
	41	+END_TILING_DATA_DEF;
	42	+
	43	+REGISTER_TILING_DATA_CLASS(MhcPostBackward, MhcPostBackwardTilingData)
	44	+
	45	+struct MhcPostBackwardCompileInfo {
	46	+};
	47	+
	48	+} // namespace optiling
	49	+
	50	+#endif // MHC_POST_BACKWARD_TILING_H

Amhc/mhc_post_backward/op_kernel/mhc_post_backward.cpp+32-0

@@ -0,0 +1,32 @@
	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 mhc_post_backward.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+#include "mhc_post_backward.h"
	17	+#include "kernel_operator.h"
	18	+
	19	+using namespace AscendC;
	20	+
	21	+extern "C" __global__ __aicore__ void mhc_post_backward(
	22	+ GM_ADDR grad_y, GM_ADDR x, GM_ADDR h_res, GM_ADDR h_out, GM_ADDR h_post,
	23	+ GM_ADDR grad_x, GM_ADDR grad_h_res, GM_ADDR grad_h_out, GM_ADDR grad_h_post,
	24	+ GM_ADDR workspace, GM_ADDR tiling)
	25	+{
	26	+ TPipe pipe;
	27	+ GET_TILING_DATA(tiling_data, tiling);
	28	+ KernelMhcPostBackward<DTYPE_GRAD_Y> op;
	29	+ op.Init(grad_y, x, h_res, h_out, h_post, grad_x, grad_h_res, grad_h_out,
	30	+ grad_h_post, tiling_data, &pipe);
	31	+ op.Process();
	32	+}

Amhc/mhc_post_backward/op_kernel/mhc_post_backward.h+438-0

@@ -0,0 +1,438 @@
	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 mhc_post_backward.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#ifndef MHC_POST_BACKWARD_H
	17	+#define MHC_POST_BACKWARD_H
	18	+
	19	+#include "kernel_operator.h"
	20	+using namespace AscendC;
	21	+
	22	+constexpr float ZERO = 0;
	23	+
	24	+template <typename T>
	25	+class KernelMhcPostBackward
	26	+{
	27	+public:
	28	+ __aicore__ inline KernelMhcPostBackward() {}
	29	+
	30	+ __aicore__ inline void Init(
	31	+ GM_ADDR grad_y, GM_ADDR x, GM_ADDR h_res, GM_ADDR h_out, GM_ADDR h_post,
	32	+ GM_ADDR grad_x, GM_ADDR grad_h_res, GM_ADDR grad_h_out, GM_ADDR grad_h_post,
	33	+ const MhcPostBackwardTilingData& tilingData, TPipe* pipe);
	34	+ __aicore__ inline void Process();
	35	+
	36	+protected:
	37	+ static constexpr uint64_t BUFFER_NUM = 1;
	38	+
	39	+ __aicore__ inline void VecMatmulMknk(
	40	+ LocalTensor<float>& A, LocalTensor<float>& B, LocalTensor<float>& C,
	41	+ LocalTensor<float>& broadcastBuffer, LocalTensor<float>& reduceBuffer, uint32_t m, uint32_t k,
	42	+ uint32_t n, uint32_t alignN);
	43	+ __aicore__ inline void VecMatmulMkkn(
	44	+ LocalTensor<float>& A, LocalTensor<float>& B, LocalTensor<float>& C,
	45	+ LocalTensor<float>& broadcastBuffer, uint32_t m, uint32_t k, uint32_t n);
	46	+
	47	+ TBuf<TPosition::VECCALC> dFPostResCastBuf, FOutCastBuf, HLPostBuf, xLCastBuf, HLResBuf;
	48	+ TBuf<TPosition::VECCALC> dHLPostBuf, dFOutCastBuf, dHLResBuf, dxLCastBuf;
	49	+ TBuf<TPosition::VECCALC> dFPostResBuf, FOutBuf, xLBuf, dFOutBuf, dxLBuf;
	50	+ TBuf<TPosition::VECCALC> dHLResTmpBuf1, dHLResTmpBuf2, dHLResTmpBuf3, dHLResTmpBuf4;
	51	+ TBuf<TPosition::VECCALC> dxLTmpBuf2;
	52	+
	53	+ GlobalTensor<T> dFPostResGm, FOutGm, xLGm;
	54	+ GlobalTensor<float> HLResGm, HLPostGm;
	55	+
	56	+ GlobalTensor<T> dxLGm, dFOutGm;
	57	+ GlobalTensor<float> dHResGm, dHPostGm;
	58	+
	59	+ LocalTensor<float> dFPostResUb, FOutUb, HLPostUb, xLUb, HLResUb;
	60	+ LocalTensor<T> dFPostResCastUb, FOutCastUb, xLCastUb, dFOutCastUb, dxLCastUb;
	61	+ LocalTensor<float> dHLPostUb, dFOutUb, dHLResUb, dxLUb;
	62	+
	63	+ LocalTensor<float> dHLResTmp1, dHLResTmp2, dHLResTmp3;
	64	+
	65	+ LocalTensor<float> dxLTmp2;
	66	+
	67	+ uint64_t coreUsed = 0;
	68	+ uint64_t singleCoreBS = 0;
	69	+ uint64_t tailBS = 0;
	70	+ uint64_t frontCore = 0;
	71	+ uint64_t tailCore = 0;
	72	+
	73	+ uint64_t blockChannel = 0;
	74	+ uint64_t channel = 0;
	75	+ uint64_t n = 0;
	76	+ uint64_t alignN = 0;
	77	+ uint64_t loopC = 0;
	78	+ uint64_t tailC = 0;
	79	+
	80	+ uint64_t dFPostResSize = 0;
	81	+ uint64_t xSize = 0;
	82	+ uint64_t hResSize = 0;
	83	+ uint64_t hOutSize = 0;
	84	+ uint64_t hPostSize = 0;
	85	+
	86	+ DataCopyPadParams padParams{false, 0, 0, 0};
	87	+};
	88	+
	89	+template <typename T>
	90	+__aicore__ inline void KernelMhcPostBackward<T>::Init(
	91	+ GM_ADDR grad_y, GM_ADDR x, GM_ADDR h_res, GM_ADDR h_out, GM_ADDR h_post,
	92	+ GM_ADDR grad_x, GM_ADDR grad_h_res, GM_ADDR grad_h_out, GM_ADDR grad_h_post,
	93	+ const MhcPostBackwardTilingData& tilingData, TPipe* pipe
	94	+)
	95	+{
	96	+ this->coreUsed = tilingData.coreUsed;
	CANN-robot4月20日代码结构与可维护性: 第94行重复赋值了this->coreUsed = tilingData.coreUsed，第98行又重复了一次。这种重复赋值没有实际意义，可能是复制粘贴错误或代码冗余，影响代码可读性。问题类型: 代码结构与可维护性文件路径: `mhc/mhc_post_backward/op_kernel/mhc_post_backward.h` 行号: 93 问题代码: this->coreUsed = tilingData.coreUsed; this->singleCoreBS = tilingData.singleCoreBS; this->tailBS = tilingData.tailBS; this->frontCore = tilingData.frontCore; this->tailCore = tilingData.tailCore; this->coreUsed = tilingData.coreUsed; 修改建议: 删除第98行的重复赋值语句，保留第94行的赋值即可。 --- 此评论由代码审查工具自动生成
	97	+ this->singleCoreBS = tilingData.singleCoreBS;
	98	+ this->tailBS = tilingData.tailBS;
	99	+ this->frontCore = tilingData.frontCore;
	100	+ this->tailCore = tilingData.tailCore;
	101	+ this->coreUsed = tilingData.coreUsed;
	ww4月23日重复赋值 fazhenyao123 4月23日评论： 4月23日评论：
	102	+
	103	+ this->blockChannel = tilingData.blockChannel;
	ww4月23日发生数据截断，tilingdata的数据类型要和kernel一致 fazhenyao123 4月25日评论： 4月25日评论：
	104	+ this->channel = tilingData.channel;
	105	+ this->n = tilingData.n;
	106	+ this->alignN = tilingData.alignN;
	107	+ this->loopC = tilingData.loopC;
	108	+ this->tailC = tilingData.tailC;
	109	+
	110	+ this->dFPostResSize = tilingData.dFPostResSize;
	111	+ this->xSize = tilingData.xSize;
	112	+ this->hResSize = tilingData.hResSize;
	113	+ this->hOutSize = tilingData.hOutSize;
	114	+ this->hPostSize = tilingData.hPostSize;
	115	+
	116	+ this->dFPostResGm.SetGlobalBuffer((__gm__ T*)grad_y, this->dFPostResSize);
	117	+ this->FOutGm.SetGlobalBuffer((__gm__ T*)h_out, this->hOutSize);
	118	+ this->xLGm.SetGlobalBuffer((__gm__ T*)x, this->xSize);
	119	+ this->HLResGm.SetGlobalBuffer((__gm__ float*)h_res, this->hResSize);
	120	+ this->HLPostGm.SetGlobalBuffer((__gm__ float*)h_post, this->hPostSize);
	121	+
	122	+ this->dxLGm.SetGlobalBuffer((__gm__ T*)grad_x, this->xSize);
	123	+ this->dFOutGm.SetGlobalBuffer((__gm__ T*)grad_h_out, this->hOutSize);
	124	+ this->dHResGm.SetGlobalBuffer((__gm__ float*)grad_h_res, this->hResSize);
	125	+ this->dHPostGm.SetGlobalBuffer((__gm__ float*)grad_h_post, this->hPostSize);
	126	+
	127	+ pipe->InitBuffer(this->dFPostResCastBuf, this->n * this->blockChannel * sizeof(T));
	128	+ pipe->InitBuffer(this->FOutCastBuf, this->blockChannel * sizeof(T));
	129	+ pipe->InitBuffer(this->HLPostBuf, this->n * sizeof(float));
	130	+ pipe->InitBuffer(this->xLCastBuf, this->n * this->blockChannel * sizeof(T));
	131	+ pipe->InitBuffer(this->HLResBuf, this->n * this->n * sizeof(float));
	132	+
	133	+ pipe->InitBuffer(this->dHLPostBuf, this->alignN * sizeof(float));
	134	+ pipe->InitBuffer(this->dFOutCastBuf, this->blockChannel * sizeof(T));
	135	+ pipe->InitBuffer(this->dHLResBuf, this->n * this->alignN * sizeof(float));
	136	+ pipe->InitBuffer(this->dxLCastBuf, this->n * this->blockChannel * sizeof(T));
	137	+
	138	+ pipe->InitBuffer(this->dFPostResBuf, this->n * this->blockChannel * sizeof(float)); // [n, tileC]
	139	+ pipe->InitBuffer(this->FOutBuf, this->blockChannel * sizeof(float)); // [1, tileC]
	140	+ pipe->InitBuffer(this->xLBuf, this->n * this->blockChannel * sizeof(float)); // [n, tileC]
	141	+ pipe->InitBuffer(this->dFOutBuf, this->blockChannel * sizeof(float)); // [1, tileC]
	142	+ pipe->InitBuffer(this->dxLBuf, this->n * this->blockChannel * sizeof(float)); // [n, tileC]
	143	+
	144	+ pipe->InitBuffer(this->dHLResTmpBuf1, this->blockChannel * this->n * sizeof(float)); // [tilC, n] [k, n]
	145	+ pipe->InitBuffer(this->dHLResTmpBuf2, this->blockChannel * this->n * this->n * sizeof(float));
	146	+ pipe->InitBuffer(this->dHLResTmpBuf3, this->n * this->alignN * sizeof(float)); // [n, n] [m, n]
	147	+ pipe->InitBuffer(this->dHLResTmpBuf4, this->n * this->n * sizeof(uint8_t)); // [n, n] [m, n]
	148	+
	149	+ pipe->InitBuffer(this->dxLTmpBuf2, this->n * this->n * sizeof(float));
	150	+}
	151	+
	152	+template <typename T>
	153	+__aicore__ inline void KernelMhcPostBackward<T>::Process()
	154	+{
	155	+ uint32_t coreId = GetBlockIdx();
	156	+
	157	+ if (coreId >= this->coreUsed) {
	158	+ return;
	159	+ }
	160	+
	161	+ uint64_t startIdx = coreId > this->frontCore ? (coreId - this->frontCore) * this->tailBS +
	162	+ this->frontCore * this->singleCoreBS : coreId * this->singleCoreBS;
	163	+
	164	+ uint64_t endIdx = startIdx + ((coreId < this->frontCore) ? this->singleCoreBS : this->tailBS);
	165	+
	166	+ this->dFPostResCastUb = this->dFPostResCastBuf.template Get<T>();
	167	+ this->FOutCastUb = this->FOutCastBuf.template Get<T>();
	168	+ this->HLPostUb = this->HLPostBuf.template Get<float>();
	169	+ this->xLCastUb = this->xLCastBuf.template Get<T>();
	170	+ this->HLResUb = this->HLResBuf.template Get<float>();
	171	+
	172	+ this->dHLPostUb = this->dHLPostBuf.template Get<float>();
	173	+ this->dFOutCastUb = this->dFOutCastBuf.template Get<T>();
	174	+ this->dHLResUb = this->dHLResBuf.template Get<float>();
	175	+ this->dxLCastUb = this->dxLCastBuf.template Get<T>();
	176	+
	177	+ this->dFPostResUb = this->dFPostResBuf.template Get<float>();
	178	+ this->FOutUb = this->FOutBuf.template Get<float>();
	179	+ this->xLUb = this->xLBuf.template Get<float>();
	180	+ this->dFOutUb = this->dFOutBuf.template Get<float>();
	181	+ this->dxLUb = this->dxLBuf.template Get<float>();
	182	+
	183	+ this->dHLResTmp1 = this->dHLResTmpBuf1.template Get<float>();
	184	+ this->dHLResTmp2 = this->dHLResTmpBuf2.template Get<float>();
	185	+ this->dHLResTmp3 = this->dHLResTmpBuf3.template Get<float>();
	186	+
	187	+ this->dxLTmp2 = this->dxLTmpBuf2.template Get<float>();
	188	+
	189	+ for (uint64_t i = startIdx; i < endIdx; i++) {
	190	+ // Hlpost 和Hlres 驻留
	191	+ SetFlag<HardEvent::V_MTE2>(0);
	192	+ WaitFlag<HardEvent::V_MTE2>(0);
	193	+
	194	+ DataCopyParams copyParamsHLPostUb{1, (uint16_t)(this->n * sizeof(float)), 0, 0};
	195	+ DataCopyPad(this->HLPostUb, this->HLPostGm[i * this->n], copyParamsHLPostUb, this->padParams);
	196	+ DataCopyParams copyParamsHLResUb{1, (uint16_t)(this->n * this->n * sizeof(float)), 0, 0};
	197	+ DataCopyPad(this->HLResUb, this->HLResGm[i * this->n * this->n], copyParamsHLResUb, this->padParams);
	198	+
	199	+ // 清0
	200	+ SetFlag<HardEvent::MTE3_V>(0);
	201	+ WaitFlag<HardEvent::MTE3_V>(0);
	202	+
	203	+ Duplicate(this->dHLPostUb, float(0.0), this->alignN);
	204	+ Duplicate(this->dHLResUb, float(0.0), this->n * this->alignN);
	205	+
	206	+ for (int j = 0; j < this->loopC; j ++) {
	207	+ uint16_t channelStride = this->channel - this->blockChannel;
	208	+ SetFlag<HardEvent::V_MTE2>(0);
	鞠鞠鞠阳4月25日 setflag/waitflag建议贴近对应的操作，否则db会开启的不好，流水有气泡 fazhenyao123 4月25日评论： 4月25日评论：
	209	+ WaitFlag<HardEvent::V_MTE2>(0);
	鞠鞠鞠阳4月25日 event_id这里不要用魔鬼数字 fazhenyao123 4月25日评论： 4月25日评论：
	210	+
	211	+ SetFlag<HardEvent::MTE3_V>(1);
	212	+ WaitFlag<HardEvent::MTE3_V>(1);
	213	+
	214	+ DataCopyParams copyParamsdFPostResUb{
	215	+ (uint16_t)this->n, (uint16_t)(this->blockChannel * sizeof(T)),
	216	+ (uint16_t)(channelStride * sizeof(T)), 0};
	217	+ DataCopyPad(
	218	+ this->dFPostResCastUb, this->dFPostResGm[i * this->n * this->channel + j * this->blockChannel],
	219	+ copyParamsdFPostResUb, this->padParams);
	220	+ DataCopyParams copyParamsFOutUb{1, (uint16_t)(this->blockChannel * sizeof(T)), 0, 0};
	221	+ DataCopyPad(
	222	+ this->FOutCastUb, this->FOutGm[i * this->channel + j * this->blockChannel],
	223	+ copyParamsFOutUb, this->padParams);
	224	+
	225	+ DataCopyParams copyParamsXLUb{
	226	+ (uint16_t)this->n, (uint16_t)(this->blockChannel * sizeof(T)),
	227	+ (uint16_t)(channelStride * sizeof(T)), 0};
	228	+ DataCopyPad(
	229	+ this->xLCastUb, this->xLGm[i * this->n * this->channel + j * this->blockChannel],
	230	+ copyParamsXLUb, this->padParams);
	231	+
	232	+ SetFlag<HardEvent::MTE2_V>(0);
	233	+ WaitFlag<HardEvent::MTE2_V>(0);
	234	+
	235	+ SetFlag<HardEvent::MTE3_V>(0);
	236	+ WaitFlag<HardEvent::MTE3_V>(0);
	237	+
	238	+ // dHLPost : Fout @ dFPostRes.T [1, blockChannel] @ [n, blockChannel]^T
	239	+ // 搬运量：[1, n, blockChannel] + [1, 1, blockChannel]
	240	+ // 输出：[1, n] 驻留在ub做累加
	241	+ Cast(
	鞠鞠鞠阳4月25日这里有多个搬入和cast，如果追求极致性能可以给这些操作独立的eventid，这样你搬入第二个的时候就可以开始算第一个了。如果非性能瓶颈可以酌情简化，保持代码。 fazhenyao123 4月25日评论： 4月25日评论：
	242	+ this->dFPostResUb, this->dFPostResCastUb, RoundMode::CAST_NONE,
	243	+ this->n * this->blockChannel); // bf16--> fp32
	244	+ Cast(this->FOutUb, this->FOutCastUb, RoundMode::CAST_NONE, this->blockChannel); // bf16--> fp32
	245	+ Cast(this->xLUb, this->xLCastUb, RoundMode::CAST_NONE, this->n * this->blockChannel); // bf16--> fp32
	246	+
	247	+ VecMatmulMknk(
	248	+ this->FOutUb, this->dFPostResUb, this->dHLPostUb, this->dHLResTmp1,
	249	+ this->dHLResTmp3, 1, this->blockChannel, this->n, this->alignN);
	250	+
	251	+ // dHres: dF@x^T //x@dF^T
	252	+ // [n, C]@[C, n] = [n, n] 驻留在ub做累加
	253	+ VecMatmulMknk(
	254	+ this->xLUb, this->dFPostResUb, this->dHLResUb, this->dHLResTmp1,
	255	+ this->dHLResTmp3, this->n, this->blockChannel, this->n, this->alignN);
	256	+
	257	+ // dFout: H_post@dF
	258	+ // [1, n]@[n, C] = [1, C] 直接搬出
	259	+ Duplicate(this->dFOutUb, float(0.0), this->blockChannel);
	260	+ VecMatmulMkkn(
	261	+ this->HLPostUb, this->dFPostResUb, this->dFOutUb, this->dHLResTmp2,
	262	+ 1, this->n, this->blockChannel);
	263	+
	264	+ // dx_l: H_res^T@dF
	265	+ // [n, n]@[n, C] = [n, C] 直接搬出,
	266	+ Duplicate(this->dxLUb, float(0.0), this->n * this->blockChannel);
	267	+ VecMatmulMkkn(
	268	+ this->HLResUb, this->dFPostResUb, this->dxLUb, this->dHLResTmp2,
	269	+ this->n, this->n, this->blockChannel);
	270	+
	271	+
	272	+ DataCopyParams copyParamsdxLGm{
	273	+ (uint16_t)this->n, (uint16_t)(this->blockChannel * sizeof(T)), 0,
	274	+ (uint16_t)(channelStride * sizeof(T))};
	275	+ DataCopyParams copyParamsdFOutGm{1, (uint16_t)(this->blockChannel * sizeof(T)), 0, 0};
	276	+ Cast(this->dFOutCastUb, this->dFOutUb, RoundMode::CAST_ROUND, this->blockChannel); // float--> bf16
	277	+ Cast(this->dxLCastUb, this->dxLUb, RoundMode::CAST_ROUND, this->n * this->blockChannel); // float--> bf16
	278	+
	279	+ SetFlag<HardEvent::V_MTE3>(0);
	280	+ WaitFlag<HardEvent::V_MTE3>(0);
	281	+ DataCopyPad(
	282	+ this->dxLGm[i * this->n * this->channel + j * this->blockChannel],
	283	+ this->dxLCastUb, copyParamsdxLGm);
	284	+ DataCopyPad(
	285	+ this->dFOutGm[i * this->channel + j * this->blockChannel],
	286	+ this->dFOutCastUb, copyParamsdFOutGm);
	287	+ }
	288	+
	289	+ if (this->tailC != 0) {
	290	+ // 尾部C的处理
	291	+ uint16_t channelStride = this->channel - this->tailC;
	292	+
	293	+ SetFlag<HardEvent::V_MTE2>(0);
	294	+ WaitFlag<HardEvent::V_MTE2>(0);
	295	+
	296	+ SetFlag<HardEvent::MTE3_V>(1);
	297	+ WaitFlag<HardEvent::MTE3_V>(1);
	298	+
	299	+ DataCopyParams copyParamsdFPostResUb{
	300	+ (uint16_t)this->n, (uint16_t)(this->tailC * sizeof(T)),
	301	+ (uint16_t)(channelStride * sizeof(T)), 0};
	302	+ DataCopyPad(
	303	+ this->dFPostResCastUb,
	304	+ this->dFPostResGm[i * this->n * this->channel + this->loopC * this->blockChannel],
	305	+ copyParamsdFPostResUb, this->padParams);
	306	+
	307	+ DataCopyParams copyParamsFOutUb{1, (uint16_t)(this->tailC * sizeof(T)), 0, 0};
	308	+ DataCopyPad(
	309	+ this->FOutCastUb,
	310	+ this->FOutGm[i * this->channel + this->loopC * this->blockChannel],
	311	+ copyParamsFOutUb, this->padParams);
	312	+ DataCopyParams copyParamsXLUb{
	313	+ (uint16_t)this->n, (uint16_t)(this->tailC * sizeof(T)),
	314	+ (uint16_t)(channelStride * sizeof(T)), 0};
	315	+ DataCopyPad(
	316	+ this->xLCastUb,
	317	+ this->xLGm[i * this->n * this->channel + this->loopC * this->blockChannel],
	318	+ copyParamsXLUb, this->padParams);
	319	+ SetFlag<HardEvent::MTE2_V>(0);
	320	+ WaitFlag<HardEvent::MTE2_V>(0);
	321	+
	322	+ SetFlag<HardEvent::MTE3_V>(0);
	323	+ WaitFlag<HardEvent::MTE3_V>(0);
	324	+
	325	+ // dHLPost : Fout @ dFPostRes.T [1, tailC] @ [n, tailC]^T
	326	+ // 搬运量：[1, n, tailC] + [1, 1, tailC]
	327	+ // 输出：[1, n] 驻留在ub做累加
	328	+ Cast(this->dFPostResUb, this->dFPostResCastUb, RoundMode::CAST_NONE, this->n * this->tailC); // bf16--> fp32
	329	+ Cast(this->FOutUb, this->FOutCastUb, RoundMode::CAST_NONE, this->tailC); // bf16--> fp32
	330	+ Cast(this->xLUb, this->xLCastUb, RoundMode::CAST_NONE, this->n * this->tailC); // bf16--> fp32
	331	+
	332	+ VecMatmulMknk(
	333	+ this->FOutUb, this->dFPostResUb, this->dHLPostUb, this->dHLResTmp1,
	334	+ this->dHLResTmp3, 1, this->tailC, this->n, this->alignN);
	335	+
	336	+ // dHres: dF@x^T
	337	+ // [n, tailC]@[tailC, n] = [n, n] 驻留在ub做累加
	338	+ VecMatmulMknk(
	339	+ this->xLUb, this->dFPostResUb, this->dHLResUb, this->dHLResTmp1,
	340	+ this->dHLResTmp3, this->n, this->tailC, this->n, this->alignN);
	341	+
	342	+ // dFout: H_post@dF
	343	+ // [1, n]@[n, tailC] = [1, tailC] 直接搬出
	344	+ Duplicate(this->dFOutUb, float(0.0), this->tailC);
	345	+ VecMatmulMkkn(
	346	+ this->HLPostUb, this->dFPostResUb, this->dFOutUb, this->dHLResTmp2,
	347	+ 1, this->n, this->tailC);
	348	+ // dx_l: H_res^T@dF
	349	+ // [n, n]@[n, tailC] = [n, tailC] 直接搬出
	350	+ Duplicate(this->dxLUb, float(0.0), this->n * this->tailC);
	351	+ VecMatmulMkkn(
	352	+ this->HLResUb, this->dFPostResUb, this->dxLUb, this->dHLResTmp2,
	353	+ this->n, this->n, this->tailC);
	354	+
	355	+ DataCopyParams copyParamsdxLGm{
	356	+ (uint16_t)this->n, (uint16_t)(this->tailC * sizeof(T)), 0,
	357	+ (uint16_t)(channelStride * sizeof(T))};
	358	+ DataCopyParams copyParamsdFOutGm{1, (uint16_t)(this->tailC * sizeof(T)), 0, 0};
	359	+ Cast(this->dFOutCastUb, this->dFOutUb, RoundMode::CAST_ROUND, this->tailC); // float--> bf16
	360	+ Cast(this->dxLCastUb, this->dxLUb, RoundMode::CAST_ROUND, this->n * this->tailC); // float--> bf16
	361	+
	362	+ SetFlag<HardEvent::V_MTE3>(0);
	363	+ WaitFlag<HardEvent::V_MTE3>(0);
	364	+ DataCopyPad(
	365	+ this->dxLGm[i * this->n * this->channel + this->loopC * this->blockChannel],
	366	+ this->dxLCastUb, copyParamsdxLGm);
	367	+ DataCopyPad(
	368	+ this->dFOutGm[i * this->channel + this->loopC * this->blockChannel],
	369	+ this->dFOutCastUb, copyParamsdFOutGm);
	370	+ }
	371	+
	372	+ DataCopyParams copyParamsdHLResGm{(uint16_t)this->n, (uint16_t)(this->n * sizeof(float)), 0, 0};
	373	+ DataCopyParams copyParamsdHLPostGm{1, (uint16_t)(this->n * sizeof(float)), 0, 0};
	374	+
	375	+ SetFlag<HardEvent::V_MTE3>(2);
	376	+ WaitFlag<HardEvent::V_MTE3>(2);
	377	+ DataCopyPad(this->dHPostGm[i * this->n], this->dHLPostUb, copyParamsdHLPostGm);
	378	+ DataCopyPad(this->dHResGm[i * this->n * this->n], this->dHLResUb, copyParamsdHLResGm);
	379	+ }
	380	+}
	381	+
	382	+template <typename T>
	383	+__aicore__ inline void KernelMhcPostBackward<T>::VecMatmulMknk(
	384	+ LocalTensor<float>& A,
	385	+ LocalTensor<float>& B,
	386	+ LocalTensor<float>& C,
	387	+ LocalTensor<float>& broadcastBuffer,
	388	+ LocalTensor<float>& reduceBuffer,
	389	+ uint32_t m, uint32_t k, uint32_t n, uint32_t alignN)
	390	+{
	391	+ // Compute matrix multiplication using vector instructions.
	392	+ uint32_t reduceShape[] = {n, k};
	393	+ uint32_t broadcastSrcShape[] = {1, k};
	394	+ constexpr bool isReuse = true;
	395	+
	396	+ auto tempBuffer = this->dHLResTmpBuf4.template Get<uint8_t>();
	397	+ for (int32_t i = 0; i < m; i++) {
	398	+
	399	+ // 逐元素相乘
	400	+ PipeBarrier<PIPE_V>();
	401	+ Mul(broadcastBuffer[0], A[i * k], B[0], k);
	402	+ Mul(broadcastBuffer[k], A[i * k], B[k], k);
	403	+ Mul(broadcastBuffer[2 * k], A[i * k], B[2 * k], k);
	404	+ Mul(broadcastBuffer[3 * k], A[i * k], B[3 * k], k);
	405	+ // 沿第0维求和
	张张张硕累4月21日【review】此处ReduceSumFP32依赖broadcastBuffer，但是没有BARR PIPE_V，会有同步问题，建议新增PIPE_V
	406	+ PipeBarrier<PIPE_V>();
	407	+ ReduceSum<float, AscendC::Pattern::Reduce::AR, isReuse>(
	408	+ reduceBuffer[i * alignN], broadcastBuffer,
	409	+ tempBuffer, reduceShape, true);
	410	+ }
	411	+
	张张张硕累4月21日【review】此处Add依赖reduceBuffer，但是没有BARR PIPE_V，会有同步问题，建议新增PIPE_V
	412	+ // 累加到输出 C
	413	+ PipeBarrier<PIPE_V>();
	414	+ Add(C, C, reduceBuffer, m * alignN);
	415	+}
	416	+
	417	+template <typename T>
	418	+__aicore__ inline void KernelMhcPostBackward<T>::VecMatmulMkkn(
	419	+ LocalTensor<float>&A,
	420	+ LocalTensor<float>&B,
	421	+ LocalTensor<float>&C,
	422	+ LocalTensor<float>&broadcastBuffer,
	423	+ uint32_t m,
	424	+ uint32_t k,
	425	+ uint32_t n
	426	+)
	427	+{
	428	+ uint32_t broadcastSrcShape[] = {m*k, 1};
	429	+ uint32_t broadcastDstShape[] = {m*k, n};
	430	+ Broadcast<float, 2, 1>(broadcastBuffer, A, broadcastDstShape, broadcastSrcShape);
	431	+ PipeBarrier<PIPE_V>();
	432	+ for (int32_t i = 0; i < m; i++) {
	433	+ for (int32_t j = 0; j < k; j++) {
	434	+ MulAddDst(C[i * n], broadcastBuffer[(im+j)n], B[j * n], n);
	435	+ }
	436	+ }
	437	+}
	CANN-robot4月20日代码结构与可维护性: 文件末尾缺少换行符，不符合代码规范。大多数代码规范要求文件以换行符结束。问题类型: 代码结构与可维护性文件路径: `mhc/mhc_post_backward/op_kernel/mhc_post_backward.h` 行号: 386 问题代码: } #endif 修改建议: 在文件末尾添加一个换行符。 --- 此评论由代码审查工具自动生成
	438	+#endif
	CANN-robot4月20日文件格式: 文件末尾（第387行）缺少换行符。根据POSIX标准，文本文件的每一行应以换行符结尾。许多工具（如`cat`、`wc`、版本控制系统如git）在处理缺少结尾换行的文件时行为可能不一致，git diff有时会标记此类文件。问题类型: 文件格式文件路径: `mhc/mhc_post_backward/op_kernel/mhc_post_backward.h` 行号: 387 问题代码: #endif 修改建议: 在文件末尾的`#endif`后添加一个换行符。 --- 此评论由代码审查工具自动生成