【社区任务】AscendC实现L2Loss算子贡献-ops-nn-AtomGit

Aexperimental/loss/l2_loss/CMakeLists.txt+20-0

@@ -0,0 +1,20 @@
	1	+# ----------------------------------------------------------------------------
	2	+# This program is free software, you can redistribute it and/or modify it.
	3	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+# This file is a part of the CANN Open Software.
	5	+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
	6	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
	8	+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	9	+# See LICENSE in the root of the software repository for the full text of the License.
	10	+# ----------------------------------------------------------------------------
	11	+
	12	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	13	+if(NOT ENABLE_TEST AND NOT BENCHMARK)
	14	+ list(REMOVE_ITEM CURRENT_DIRS tests)
	15	+endif()
	16	+foreach(SUB_DIR ${CURRENT_DIRS})
	17	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	18	+ add_subdirectory(${SUB_DIR})
	19	+ endif()
	20	+endforeach()

Aexperimental/loss/l2_loss/README.md+35-0

@@ -0,0 +1,35 @@
	1	+# L2Loss
	2	+
	3	+## 产品支持情况
	4	+
	5	+\| 产品 \| 是否支持 \|
	6	+\| :--- \| :------- \|
	7	+\| Atlas A2 训练系列产品/Atlas A2 推理系列产品 \| √ \|
	8	+
	9	+## 功能说明
	10	+
	11	+- 算子功能：对输入张量 x 的所有元素计算 L2 Loss，即 `output = sum(x * x) / 2`，返回一个标量。
	12	+
	13	+## 参数说明
	14	+
	15	+\| 参数名 \| 输入/输出 \| 描述 \| 数据类型 \| 数据格式 \|
	16	+\| :----- \| :-------- \| :--- \| :------- \| :------- \|
	17	+\| x \| 输入 \| 输入张量 \| float32、float16、bfloat16 \| ND \|
	18	+\| y \| 输出 \| 输出标量（shape 为 [1]） \| float32、float16、bfloat16 \| ND \|
	19	+
	20	+## 约束说明
	21	+
	22	+- 本算子验证基于 CANN 9.0.0。
	23	+- 输入张量 x 不能为空张量（元素数量必须大于 0）。
	24	+
	25	+## 调用说明
	26	+
	27	+\| 调用方式 \| 样例代码 \| 说明 \|
	28	+\| :------- \| :------- \| :--- \|
	29	+\| aclnn 接口 \| [test_aclnn_l2_loss](examples/test_aclnn_l2_loss.cpp) \| 通过 aclnnL2Loss 接口调用 l2_loss 算子。 \|
	30	+
	31	+## 贡献说明
	32	+
	33	+本算子由社区开发者贡献，贡献流程请参考 [CANN 算子贡献指南](https://gitcode.com/cann/ops-math/blob/master/CONTRIBUTING.md)。
	34	+
	35	+欢迎提交 Issue 或 Pull Request 参与共建。如有问题，请在对应仓库的 Issue 区描述复现步骤和环境信息。

Aexperimental/loss/l2_loss/examples/test_aclnn_l2_loss.cpp+127-0

@@ -0,0 +1,127 @@
	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	+#include <iostream>
	11	+#include <vector>
	12	+#include "acl/acl.h"
	13	+#include "aclnn_l2_loss.h"
	14	+
	15	+#define CHECK_RET(cond, return_expr) \
	16	+ do { \
	17	+ if (!(cond)) { \
	18	+ printf("[CHECK_RET FAILED] %s:%d\n", __FILE__, __LINE__); \
	19	+ return_expr; \
	20	+ } \
	21	+ } while (0)
	22	+
	23	+#define LOG_PRINT(message, ...) \
	24	+ do { \
	25	+ printf(message, ##__VA_ARGS__); \
	26	+ } while (0)
	27	+
	28	+int64_t GetShapeSize(const std::vector<int64_t>& shape)
	29	+{
	30	+ int64_t shapeSize = 1;
	31	+ for (auto i : shape) {
	32	+ shapeSize *= i;
	33	+ }
	34	+ return shapeSize;
	35	+}
	36	+
	37	+
	38	+
	39	+int Init(int32_t deviceId, aclrtStream* stream)
	40	+{
	41	+ auto ret = aclInit(nullptr);
	42	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
	43	+ ret = aclrtSetDevice(deviceId);
	44	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
	45	+ ret = aclrtCreateStream(stream);
	46	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
	47	+ return 0;
	48	+}
	49	+
	50	+template <typename T>
	51	+int CreateAclTensor(
	52	+ const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType,
	53	+ aclTensor** tensor)
	54	+{
	55	+ auto size = GetShapeSize(shape) * sizeof(T);
	56	+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
	57	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
	58	+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
	59	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret);
	60	+
	61	+ std::vector<int64_t> strides(shape.size(), 1);
	62	+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
	63	+ strides[i] = shape[i + 1] * strides[i + 1];
	64	+ }
	65	+ *tensor = aclCreateTensor(
	66	+ shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(),
	67	+ *deviceAddr);
	68	+ return 0;
	69	+}
	70	+int main()
	71	+{
	72	+ int32_t deviceId = 0;
	73	+ aclrtStream stream;
	74	+ auto ret = Init(deviceId, &stream);
	75	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
	76	+
	77	+ aclTensor* selfX = nullptr;
	78	+ void* selfXDeviceAddr = nullptr;
	79	+ std::vector<int64_t> selfXShape = {2,8};
	80	+ std::vector<float> selfXHostData(16, 1.0);
	81	+ ret = CreateAclTensor(selfXHostData, selfXShape, &selfXDeviceAddr, aclDataType::ACL_FLOAT, &selfX);
	82	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	83	+
	84	+ aclTensor* out = nullptr;
	85	+ void* outDeviceAddr = nullptr;
	86	+ // L2Loss 输出为标量，shape = {1}
	87	+ std::vector<int64_t> outShape = {1};
	88	+ std::vector<float> outHostData(1, 0.0);
	89	+ ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out);
	90	+ CHECK_RET(ret == ACL_SUCCESS, return ret);
	91	+
	92	+ uint64_t workspaceSize = 0;
	93	+ aclOpExecutor* executor;
	94	+
	95	+ // L2Loss op_def 中无 attribute，接口只有 x（输入）和 out（输出）
	96	+ ret = aclnnL2LossGetWorkspaceSize(selfX, out, &workspaceSize, &executor);
	97	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnL2LossGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
	98	+ void* workspaceAddr = nullptr;
	99	+ if (workspaceSize > static_cast<uint64_t>(0)) {
	100	+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
	101	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
	102	+ }
	103	+ ret = aclnnL2Loss(workspaceAddr, workspaceSize, executor, stream);
	104	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnL2Loss failed. ERROR: %d\n", ret); return ret);
	105	+ ret = aclrtSynchronizeStream(stream);
	106	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
	107	+ // 将结果从 device 拷回 host 并打印（与 sqrt example 写法一致）
	108	+ auto size = GetShapeSize(outShape);
	109	+ std::vector<float> resultData(size, 0);
	110	+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]),
	111	+ outDeviceAddr, size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST);
	112	+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
	113	+ for (int64_t i = 0; i < size; i++) {
	114	+ LOG_PRINT("aclnnL2Loss result[%ld] is: %f\n", i, resultData[i]);
	115	+ }
	116	+ aclDestroyTensor(selfX);
	117	+ aclDestroyTensor(out);
	118	+ aclrtFree(selfXDeviceAddr);
	119	+ aclrtFree(outDeviceAddr);
	120	+ if (workspaceSize > static_cast<uint64_t>(0)) {
	121	+ aclrtFree(workspaceAddr);
	122	+ }
	123	+ aclrtDestroyStream(stream);
	124	+ aclrtResetDevice(deviceId);
	125	+ aclFinalize();
	126	+ return 0;
	127	+}

Aexperimental/loss/l2_loss/op_host/CMakeLists.txt+12-0

@@ -0,0 +1,12 @@
	1	+# ----------------------------------------------------------------------------
	2	+# This program is free software, you can redistribute it and/or modify it.
	3	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+# This file is a part of the CANN Open Software.
	5	+# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
	6	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	7	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
	8	+# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	9	+# See LICENSE in the root of the software repository for the full text of the License.
	10	+# ----------------------------------------------------------------------------
	11	+
	12	+add_modules_sources(HOSTNAME ${OPHOST_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR} OPTYPE l2_loss ACLNNTYPE aclnn)

Aexperimental/loss/l2_loss/op_host/l2_loss_def.cpp+37-0

@@ -0,0 +1,37 @@
	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 l2_loss_def.cpp
	12	+ * \brief
	13	+ */
	14	+#include "register/op_def_registry.h"
	15	+
	16	+namespace ops {
	17	+class L2Loss : public OpDef {
	18	+public:
	19	+ explicit L2Loss(const char* name) : OpDef(name)
	20	+ {
	21	+ this->Input("x")
	22	+ .ParamType(REQUIRED)
	23	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
	24	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	25	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	26	+ .AutoContiguous();
	27	+ this->Output("y")
	28	+ .ParamType(REQUIRED)
	29	+ .DataType({ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16})
	30	+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	31	+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
	32	+ .AutoContiguous();
	33	+ this->AICore().AddConfig("ascend910b");
	34	+ }
	35	+};
	36	+OP_ADD(L2Loss);
	37	+} // namespace ops

Aexperimental/loss/l2_loss/op_host/l2_loss_infershape.cpp+43-0

@@ -0,0 +1,43 @@
	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 l2_loss_infershape.cpp
	12	+ * \brief
	13	+ */
	14	+#include "register/op_impl_registry.h"
	15	+#include "log/log.h"
	16	+
	17	+using namespace ge;
	18	+
	19	+namespace ops {
	20	+static constexpr int64_t IDX_0 = 0;
	21	+
	22	+static ge::graphStatus InferShapeL2Loss(gert::InferShapeContext* context)
	23	+{
	24	+ OP_LOGD(context->GetNodeName(), "Begin to do InferShapeL2Loss");
	25	+
	26	+ // get input shapes
	27	+ const gert::Shape* xShape = context->GetInputShape(IDX_0);
	28	+ OP_CHECK_NULL_WITH_CONTEXT(context, xShape);
	29	+
	30	+ // get output shapes
	31	+ gert::Shape* yShape = context->GetOutputShape(IDX_0);
	32	+ OP_CHECK_NULL_WITH_CONTEXT(context, yShape);
	33	+
	34	+ // L2Loss output is a scalar: sum(x * x) / 2
	35	+ yShape->SetDimNum(1);
	36	+ yShape->SetDim(0, 1);
	37	+
	38	+ OP_LOGD(context->GetNodeName(), "End to do InferShapeL2Loss");
	39	+ return GRAPH_SUCCESS;
	40	+}
	41	+
	42	+IMPL_OP_INFERSHAPE(L2Loss).InferShape(InferShapeL2Loss);
	43	+} // namespace ops

Aexperimental/loss/l2_loss/op_host/l2_loss_tiling.cpp+235-0

@@ -0,0 +1,235 @@
	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 l2_loss_tiling.cpp
	12	+ * \brief
	13	+ */
	14	+#include "log/log.h"
	15	+#include "util/math_util.h"
	16	+#include "util/platform_util.h"
	17	+#include "op_host/tiling_util.h"
	18	+#include "tiling/platform/platform_ascendc.h"
	19	+#include "register/op_impl_registry.h"
	20	+#include "op_host/tiling_templates_registry.h"
	21	+#include "../op_kernel/l2_loss_tiling_data.h"
	22	+#include "../op_kernel/l2_loss_tiling_key.h"
	23	+
	24	+#include <array>
	25	+
	26	+
	27	+namespace optiling {
	28	+
	29	+ using namespace Ops::NN::OpTiling;
	30	+ constexpr uint32_t BUFFER_NUM = 2;
	31	+ constexpr uint32_t WS_SYS_SIZE = 16U * 1024U * 1024U; // 16MB
	32	+ // 与 kernel 中保持一致：固定 UB buffer 大小（不随 tile 规模变化）
	33	+ constexpr uint64_t K_DATA_CACHE_CLEAN_NEED = 64; // cache line bytes
	34	+ constexpr uint64_t K_SLOT_STRIDE = K_DATA_CACHE_CLEAN_NEED / sizeof(float); // = 16 floats
	35	+ // tmpBuffer(16 floats) + tileSumBuf(16 floats)
	36	+ constexpr uint64_t FIXED_UB_BYTES = 2 * K_SLOT_STRIDE * sizeof(float); // 128 B
	37	+ struct L2LossCompileInfo {};
	38	+ struct L2LossShapeInfo {
	39	+ uint64_t inputNum{0};
	40	+ uint64_t inputBytes{0};
	41	+ uint64_t tileBlockNum{0};
	42	+ uint64_t tileDataNum{0};
	43	+ uint64_t inputLengthAlign32{0};
	44	+ uint64_t smallCoreDataNum{0};
	45	+ uint64_t bigCoreDataNum{0};
	46	+ uint64_t smallTailDataNum{0};
	47	+ uint64_t bigTailDataNum{0};
	48	+ uint64_t finalSmallTileNum{0};
	49	+ uint64_t finalBigTileNum{0};
	50	+ uint64_t tailBlockNum{0};
	51	+ uint32_t blockSize{0};
	52	+
	53	+ };
	54	+
	55	+ static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum)
	56	+ {
	57	+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
	58	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
	59	+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
	60	+ coreNum = ascendcPlatform.GetCoreNumAiv();
	61	+ if (coreNum == 0) {
	62	+ coreNum = ascendcPlatform.GetCoreNum();
	63	+ }
	64	+ OP_CHECK_IF(coreNum <= 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED);
	65	+ OP_CHECK_IF(ubSize <= 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED);
	66	+ return ge::GRAPH_SUCCESS;
	67	+ }
	68	+
	69	+ static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context)
	70	+ {
	71	+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
	72	+ size_t usrSize = WS_SYS_SIZE;
	73	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
	74	+ uint32_t sysWorkspaceSize = ascendcPlatform.GetLibApiWorkSpaceSize();
	75	+ size_t* currentWorkspace = context->GetWorkspaceSizes(1);
	76	+ currentWorkspace[0] = usrSize + sysWorkspaceSize;
	77	+ return ge::GRAPH_SUCCESS;
	78	+ }
	79	+
	80	+ static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, uint64_t ubSize, L2LossShapeInfo& info)
	81	+ {
	82	+ OP_CHECK_IF(
	83	+ context == nullptr \|\| context->GetInputShape(0) == nullptr \|\| context->GetInputDesc(0) == nullptr, OP_LOGE(context, "context or input is nullptr"),
	84	+ return ge::GRAPH_FAILED);
	85	+ int64_t shapeSize = context->GetInputShape(0)->GetStorageShape().GetShapeSize();
	86	+ if (shapeSize <= 0) {
	87	+ OP_LOGE(context, "inputNum is non-positive (%ld), shape may contain unknown dimensions", (long)shapeSize);
	88	+ return ge::GRAPH_FAILED;
	89	+ }
	90	+ info.inputNum = static_cast<uint64_t>(shapeSize);
	91	+
	92	+ uint32_t typeLength = 0;
	93	+ ge::TypeUtils::GetDataTypeLength(context->GetInputDesc(0)->GetDataType(), typeLength);
	94	+ info.inputBytes = typeLength;
	95	+ if (info.inputBytes == 0) {
	96	+ OP_LOGE(context, "inputBytes is 0, invalid number");
	97	+ return ge::GRAPH_FAILED;
	98	+ }
	99	+
	100	+ uint64_t inputLength = info.inputNum * info.inputBytes;
	101	+ // 扣除固定 buffer 后的可用 UB（tmpBuffer + tileSumBuf 共 128B）
	102	+ uint64_t effectiveUbSize = (ubSize > FIXED_UB_BYTES) ? (ubSize - FIXED_UB_BYTES) : 0;
	103	+ // per-element UB：双缓冲输入队列 + half/bf16 专用的 Cast 上行 buffer（tmpFloat）
	104	+ // float: 2×4 + 0 = 8 B/elem
	105	+ // half/bf16: 2×2 + 4 = 8 B/elem （两者恰好相同）
	106	+ uint64_t tmpFloatPerElem = (info.inputBytes == sizeof(float)) ? 0ULL : sizeof(float);
	107	+ uint64_t ubBytesPerElem = BUFFER_NUM * info.inputBytes + tmpFloatPerElem;
	108	+ info.tileDataNum = effectiveUbSize / ubBytesPerElem;
	109	+ // 对齐到 blockSize/inputBytes 的整数倍（DMA 要求 blockSize 对齐）
	110	+ uint64_t elemsPerBlock = info.blockSize / info.inputBytes;
	111	+ info.tileDataNum = (info.tileDataNum / elemsPerBlock) * elemsPerBlock;
	112	+ info.tileBlockNum = info.tileDataNum * info.inputBytes / info.blockSize;
	113	+ info.inputLengthAlign32 = (((inputLength + info.blockSize - 1) / info.blockSize) * info.blockSize);
	114	+ return ge::GRAPH_SUCCESS;
	115	+ }
	116	+
	117	+ static ge::graphStatus CalculateCoreBlockNums(int64_t coreNum, L2LossShapeInfo& info)
	118	+ {
	119	+ if(0 == coreNum \|\| 0 == info.tileBlockNum) {
	120	+ return ge::GRAPH_FAILED;
	121	+ }
	122	+ uint64_t everyCoreInputBlockNum = info.inputLengthAlign32 / info.blockSize / coreNum;
	123	+ info.tailBlockNum = (info.inputLengthAlign32 / info.blockSize) % coreNum;
	124	+ info.smallCoreDataNum = everyCoreInputBlockNum * info.blockSize / info.inputBytes;
	125	+ uint64_t smallTileNum = everyCoreInputBlockNum / info.tileBlockNum;
	126	+ info.finalSmallTileNum = (everyCoreInputBlockNum % info.tileBlockNum) == 0 ? smallTileNum : smallTileNum + 1;
	127	+ info.smallTailDataNum = info.smallCoreDataNum - (info.tileDataNum * smallTileNum);
	128	+ info.smallTailDataNum = info.smallTailDataNum == 0 ? info.tileDataNum : info.smallTailDataNum;
	129	+
	130	+ everyCoreInputBlockNum += 1;
	131	+ info.bigCoreDataNum = everyCoreInputBlockNum * info.blockSize / info.inputBytes;
	132	+ uint64_t bigTileNum = everyCoreInputBlockNum / info.tileBlockNum;
	133	+ info.finalBigTileNum = (everyCoreInputBlockNum % info.tileBlockNum) == 0 ? bigTileNum : bigTileNum + 1;
	134	+ info.bigTailDataNum = info.bigCoreDataNum - info.tileDataNum * bigTileNum;
	135	+ info.bigTailDataNum = info.bigTailDataNum == 0 ? info.tileDataNum : info.bigTailDataNum;
	136	+ // bigTailDataNum/smallTailDataNum 是两个 BLOCK_SIZE/inputBytes 整数倍之差，自然对齐，无需 round-up
	137	+
	138	+ return ge::GRAPH_SUCCESS;
	139	+ }
	140	+
	141	+ // 根据元素数量直接查表得到最优启动核数。
	142	+ static int64_t PickCoresFromElems(uint64_t numElems, int64_t maxAvailableCores, bool isFloat32) {
	143	+ struct Bucket { uint64_t upperElems; int64_t cores; };
	144	+ constexpr uint64_t K = 1024ULL;
	145	+ constexpr uint64_t M = 1024ULL * K;
	146	+ // fp32 扫描结果：32K→2, 1M→12, 2M→16, 4M→24, 8M→28, 32M→32
	147	+ static constexpr std::array<Bucket, 7> kTableFp32 = {{
	148	+ { 32ULL * K, 2 },
	149	+ { 1ULL * M, 12 },
	150	+ { 2ULL * M, 16 },
	151	+ { 4ULL * M, 24 },
	152	+ { 8ULL * M, 28 },
	153	+ { 32ULL * M, 32 },
	154	+ { UINT64_MAX, 32 },
	155	+ }};
	156	+ // fp16 扫描结果：32K→4, 1M→12, 4M→20, 8M→24, 16M→32, 32M→40
	157	+ static constexpr std::array<Bucket, 7> kTableFp16 = {{
	158	+ { 32ULL * K, 4 },
	159	+ { 1ULL * M, 12 },
	160	+ { 4ULL * M, 20 },
	161	+ { 8ULL * M, 24 },
	162	+ { 16ULL * M, 32 },
	163	+ { 32ULL * M, 40 },
	164	+ { UINT64_MAX, 40 },
	165	+ }};
	166	+ const auto& table = isFloat32 ? kTableFp32 : kTableFp16;
	167	+ for (const auto& b : table) {
	168	+ if (numElems <= b.upperElems) {
	169	+ return std::min<int64_t>(b.cores, maxAvailableCores);
	170	+ }
	171	+ }
	172	+ return maxAvailableCores;
	173	+ }
	174	+
	175	+
	176	+ // tiling 分发入口
	177	+ static ge::graphStatus L2LossTilingFunc(gert::TilingContext* context)
	178	+ {
	179	+ OP_CHECK_IF(context == nullptr, OP_LOGE(context, "context is nullptr"), return ge::GRAPH_FAILED);
	180	+ L2LossTilingData* tiling = context->GetTilingData<L2LossTilingData>();
	181	+ L2LossShapeInfo shapeInfo;
	182	+ OP_CHECK_NULL_WITH_CONTEXT(context, tiling);
	183	+ OP_CHECK_IF(
	184	+ memset_s(tiling, sizeof(L2LossTilingData), 0, sizeof(L2LossTilingData)) != EOK,
	185	+ OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED);
	186	+ shapeInfo.blockSize = Ops::Base::GetUbBlockSize(context);
	187	+ uint64_t ubSize;
	188	+ int64_t coreNum;
	189	+ ge::graphStatus ret = GetPlatformInfo(context, ubSize, coreNum);
	190	+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED);
	191	+ ret = GetShapeAttrsInfo(context, ubSize, shapeInfo);
	192	+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED);
	193	+ if (shapeInfo.tileDataNum >= shapeInfo.inputNum) {
	194	+ coreNum = 1;
	195	+ } else {
	196	+ int64_t maxCoresByBlocks = static_cast<int64_t>(shapeInfo.inputLengthAlign32 / shapeInfo.blockSize);
	197	+ // 按元素数查表限核，fp32 和 fp16/bf16 各有独立拐点。
	198	+ bool isFloat32 = (shapeInfo.inputBytes == sizeof(float));
	199	+ int64_t requested = PickCoresFromElems(shapeInfo.inputNum, coreNum, isFloat32);
	200	+ coreNum = std::max<int64_t>(1, std::min({coreNum, maxCoresByBlocks, requested}));
	201	+ }
	202	+ ret = CalculateCoreBlockNums(coreNum, shapeInfo);
	203	+ if (ret != ge::GRAPH_SUCCESS) {
	204	+ OP_LOGE(context, "coreNum or tileBlockNum is 0, invalid number");
	205	+ return ret;
	206	+ }
	207	+ // 溢出保护：确保关键字段不超出 uint32_t 范围
	208	+ OP_CHECK_IF(shapeInfo.bigCoreDataNum > UINT32_MAX \|\| shapeInfo.smallCoreDataNum > UINT32_MAX \|\|
	209	+ shapeInfo.tileDataNum > UINT32_MAX \|\| shapeInfo.inputNum > UINT32_MAX,
	210	+ OP_LOGE(context, "tiling data overflow: shape too large for uint32_t"),
	211	+ return ge::GRAPH_FAILED);
	212	+ tiling->smallCoreDataNum = static_cast<uint32_t>(shapeInfo.smallCoreDataNum);
	213	+ tiling->bigCoreDataNum = static_cast<uint32_t>(shapeInfo.bigCoreDataNum);
	214	+ tiling->tileDataNum = static_cast<uint32_t>(shapeInfo.tileDataNum);
	215	+ tiling->smallTailDataNum = static_cast<uint32_t>(shapeInfo.smallTailDataNum);
	216	+ tiling->bigTailDataNum = static_cast<uint32_t>(shapeInfo.bigTailDataNum);
	217	+ tiling->finalSmallTileNum = static_cast<uint32_t>(shapeInfo.finalSmallTileNum);
	218	+ tiling->finalBigTileNum = static_cast<uint32_t>(shapeInfo.finalBigTileNum);
	219	+ tiling->tailBlockNum = static_cast<uint32_t>(shapeInfo.tailBlockNum);
	220	+ tiling->inputNum = static_cast<uint32_t>(shapeInfo.inputNum);
	221	+ tiling->blockNum = static_cast<uint32_t>(coreNum);
	222	+ OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"), return ge::GRAPH_FAILED);
	223	+ context->SetTilingKey(coreNum == 1 ? ELEMENTWISE_TPL_SCH_MODE_0 : ELEMENTWISE_TPL_SCH_MODE_1);
	224	+ context->SetBlockDim(coreNum);
	225	+ return ge::GRAPH_SUCCESS;
	226	+ }
	227	+
	228	+ static ge::graphStatus TilingParseForL2Loss([[maybe_unused]] gert::TilingParseContext* context)
	229	+ {
	230	+ return ge::GRAPH_SUCCESS;
	231	+ }
	232	+
	233	+ // tiling注册入口.
	234	+ IMPL_OP_OPTILING(L2Loss).Tiling(L2LossTilingFunc).TilingParse<L2LossCompileInfo>(TilingParseForL2Loss);
	235	+} // namespace optiling

Aexperimental/loss/l2_loss/op_kernel/l2_loss.cpp+24-0

@@ -0,0 +1,24 @@
	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 l2_loss.cpp
	12	+ * \brief
	13	+*/
	14	+#include "l2_loss.h"
	15	+
	16	+template <uint32_t schMode>
	17	+__global__ __aicore__ void l2_loss(GM_ADDR x, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling)
	18	+{
	19	+ REGISTER_TILING_DEFAULT(L2LossTilingData);
	20	+ GET_TILING_DATA_WITH_STRUCT(L2LossTilingData, tilingData, tiling);
	21	+ NsL2Loss::L2Loss<DTYPE_X, schMode> op;
	22	+ op.Init(x, z, workspace, &tilingData);
	23	+ op.Process();
	24	+}

Aexperimental/loss/l2_loss/op_kernel/l2_loss.h+269-0

@@ -0,0 +1,269 @@
	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 l2_loss.h
	12	+ * \brief
	13	+ */
	14	+#ifndef L2_LOSS_H
	15	+#define L2_LOSS_H
	16	+
	17	+#include "kernel_operator.h"
	18	+#include "kernel_tiling/kernel_tiling.h"
	19	+#include "l2_loss_tiling_data.h"
	20	+#include "l2_loss_tiling_key.h"
	21	+
	22	+namespace NsL2Loss {
	23	+
	24	+using namespace AscendC;
	25	+
	26	+constexpr int32_t BUFFER_NUM = 2;
	27	+constexpr int32_t DATA_CACHE_CLEAN_NEED = 64;
	28	+constexpr int32_t SLOT_STRIDE = DATA_CACHE_CLEAN_NEED / sizeof(float);
	29	+constexpr float INV_SQRT2 = 0.70710678118654752f; // 1/√2 这是规定做法，不要考虑改这个
	30	+template <typename T, uint32_t schMode>
	31	+class L2Loss {
	32	+public:
	33	+ __aicore__ inline L2Loss(){};
	34	+
	35	+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, GM_ADDR workspace, const L2LossTilingData* tilingData);
	36	+ __aicore__ inline void Process();
	37	+
	38	+private:
	39	+ __aicore__ inline void CopyIn(int32_t progress);
	40	+ __aicore__ inline void L2LossAxesAll();
	41	+
	42	+private:
	43	+ AscendC::TPipe pipe;
	44	+ AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueInput;
	45	+
	46	+ AscendC::GlobalTensor<T> xGm;
	47	+ AscendC::GlobalTensor<T> zGm;
	48	+ AscendC::GlobalTensor<float> workGm;
	49	+
	50	+ // tileDataNum * sizeof(float) bytes
	51	+ // 仅 T != float 时真正用到：作为 Cast 目标（将输入从 T 升精度为 float）
	52	+ // T == float 时该 buffer 未被实际读写——Cast 跳过，ReduceSum 的 sharedTmpBuffer
	53	+ AscendC::TBuf<AscendC::TPosition::VECCALC> tmpFloat;
	54	+
	55	+ // SLOT_STRIDE * sizeof(float) = 16 * 4 = 64 bytes（一条 Cache Line）
	56	+ // 用于原子写回前的暂存：将 ReduceSum 结果写入此 buf，再经 SetAtomicAdd 搬至 workGm
	57	+ AscendC::TBuf<AscendC::TPosition::VECCALC> tmpBuffer;
	58	+
	59	+ // SLOT_STRIDE * sizeof(float) = 16 * 4 = 64 bytes（32B 对齐，满足 ReduceSum dst 要求）
	60	+ // ReduceSum 的输出目标，存放当前核对 accum 向量的标量归约结果
	61	+ AscendC::TBuf<AscendC::TPosition::VECCALC> tileSumBuf;
	62	+
	63	+ uint32_t blockIdx;
	於於於欣洁6月4日 `globalBufferIndex`、`coreDataNum`、`tileDataNum` 等均为 `uint32_t`，`tilingData->bigCoreDataNum * blockIdx` 和 `progress * tileDataNum` 都在 32 位范围内计算。与 host 侧截断问题叠加后，大输入场景可能 GM offset 回绕。
	64	+ uint32_t blockNum;
	65	+ uint64_t globalOffset;
	66	+
	67	+ uint32_t coreDataNum;
	68	+ uint32_t tileNum;
	69	+ uint32_t tileDataNum;
	70	+ uint32_t tailDataNum;
	71	+ uint32_t validCoreDataNum;
	72	+ uint32_t lastTileValidLen;
	73	+};
	74	+
	75	+template <typename T, uint32_t schMode>
	76	+__aicore__ inline void L2Loss<T, schMode>::Init(GM_ADDR x, GM_ADDR z, GM_ADDR workspace, const L2LossTilingData* tilingData)
	77	+{
	78	+ ASSERT(AscendC::GetBlockNum() != 0 && "block dim can not be zero!");
	79	+ this->blockIdx = AscendC::GetBlockIdx();
	80	+ this->blockNum = AscendC::GetBlockNum();
	81	+ uint64_t globalBufferIndex = static_cast<uint64_t>(tilingData->bigCoreDataNum) * this->blockIdx;
	82	+ this->tileDataNum = tilingData->tileDataNum;
	83	+
	84	+ if (this->blockIdx < tilingData->tailBlockNum) {
	85	+ this->coreDataNum = tilingData->bigCoreDataNum;
	86	+ this->tileNum = tilingData->finalBigTileNum;
	87	+ this->tailDataNum = tilingData->bigTailDataNum;
	88	+ } else {
	89	+ this->coreDataNum = tilingData->smallCoreDataNum;
	90	+ this->tileNum = tilingData->finalSmallTileNum;
	91	+ this->tailDataNum = tilingData->smallTailDataNum;
	92	+ globalBufferIndex -= static_cast<uint64_t>(tilingData->bigCoreDataNum - tilingData->smallCoreDataNum) * (this->blockIdx - tilingData->tailBlockNum);
	93	+ }
	94	+
	95	+ // Compute the valid data count for this core to exclude padding added at tail.
	96	+ uint64_t inputNum64 = static_cast<uint64_t>(tilingData->inputNum);
	97	+ this->validCoreDataNum = inputNum64 > globalBufferIndex ?
	98	+ static_cast<uint32_t>(inputNum64 - globalBufferIndex) : 0;
	99	+ if (this->validCoreDataNum > this->coreDataNum) {
	100	+ this->validCoreDataNum = this->coreDataNum;
	101	+ }
	102	+
	103	+ // tiling 层已保证 coreDataNum/tailDataNum 是 BLOCK_SIZE/inputBytes 的整数倍，
	104	+ // 内核可直接使用，无需重算或裁剪。
	105	+ xGm.SetGlobalBuffer((__gm__ T*)x + globalBufferIndex, this->coreDataNum);
	106	+ zGm.SetGlobalBuffer((__gm__ T*)z, 1);
	107	+ // float 单核：workGm 直接指向 z（仅用来承接结果，实际写入走 zGm）
	108	+ // float 多核：用 workspace 做原子累加中转，最后由 block 0 写回 z
	109	+ // half/bf16：始终用 workspace 做 float 中转
	110	+ uint32_t workGmSize = SLOT_STRIDE;
	111	+ if constexpr (std::is_same_v<T, float> && schMode == ELEMENTWISE_TPL_SCH_MODE_0) {
	112	+ workGm.SetGlobalBuffer((__gm__ float*)z, workGmSize);
	113	+ } else {
	114	+ workGm.SetGlobalBuffer((__gm__ float*)workspace, workGmSize);
	115	+ }
	116	+ if constexpr (schMode != ELEMENTWISE_TPL_SCH_MODE_0) {
	117	+ if (AscendC::GetBlockIdx() == 0) {
	118	+ AscendC::InitGlobalMemory(workGm, workGmSize, (float)0.0f);
	119	+ }
	120	+ AscendC::SyncAll<true>();
	121	+ }
	122	+ pipe.InitBuffer(inQueueInput, BUFFER_NUM, this->tileDataNum * sizeof(T));
	123	+ // float 输入直接原地操作（ReinterpretCast），无需 Cast 中转 buffer
	124	+ // half/bf16 需要将输入 Cast 升精度到 float，才需要 tmpFloat
	125	+ if constexpr (!std::is_same_v<T, float>) {
	126	+ pipe.InitBuffer(tmpFloat, tileDataNum * sizeof(float));
	127	+ }
	128	+ pipe.InitBuffer(tmpBuffer, workGmSize * sizeof(float));
	129	+ pipe.InitBuffer(tileSumBuf, SLOT_STRIDE * sizeof(float)); // ReduceSum output, 32B aligned
	130	+
	131	+ // 预计算最后一个 tile 的有效长度，供 CopyIn 使用 DataCopyPad
	132	+ uint64_t processedBeforeLastTile = static_cast<uint64_t>(this->tileNum - 1) * this->tileDataNum;
	133	+ this->lastTileValidLen = (this->validCoreDataNum > processedBeforeLastTile)
	134	+ ? static_cast<uint32_t>(this->validCoreDataNum - processedBeforeLastTile)
	135	+ : 0;
	136	+
	137	+ this->globalOffset = globalBufferIndex;
	138	+}
	139	+
	140	+template <typename T, uint32_t schMode>
	141	+__aicore__ inline void L2Loss<T, schMode>::CopyIn(int32_t progress)
	142	+{
	143	+ uint64_t gmOffset = static_cast<uint64_t>(progress) * this->tileDataNum;
	144	+ AscendC::LocalTensor<T> xLocal = inQueueInput.AllocTensor<T>();
	145	+ if (progress == (int32_t)(tileNum - 1)) {
	146	+ // 最后一个 tile 使用 DataCopyPad，按有效元素字节数搬运，硬件自动 pad 零
	147	+ AscendC::DataCopyExtParams copyParams{1, static_cast<uint32_t>(this->lastTileValidLen * sizeof(T)), 0, 0, 0};
	148	+ AscendC::DataCopyPadExtParams<T> padParams{true, 0, 0, static_cast<T>(0)};
	149	+ AscendC::DataCopyPad(xLocal, xGm[gmOffset], copyParams, padParams);
	150	+ } else {
	151	+ AscendC::DataCopy(xLocal, xGm[gmOffset], tileDataNum);
	152	+ }
	153	+ inQueueInput.EnQue(xLocal);
	154	+}
	155	+template <typename T, uint32_t schMode>
	156	+__aicore__ inline void L2Loss<T, schMode>::Process()
	157	+{
	158	+ L2LossAxesAll();
	159	+}
	160	+
	161	+template <typename T, uint32_t schMode>
	162	+__aicore__ inline void L2Loss<T, schMode>::L2LossAxesAll()
	163	+{
	164	+ const uint32_t loopCount = this->tileNum;
	165	+ const uint32_t tileLen = this->tileDataNum;
	166	+ const uint32_t lastTileLen = this->tailDataNum;
	167	+
	168	+ // localSum 作为寄存器变量，不占 UB
	169	+ float localSum = 0.0f;
	170	+
	171	+ // 循环外提前获取固定 buffer 地址，避免每次循环重复计算 UB 偏移
	172	+ AscendC::LocalTensor<float> tileSumTensor = tileSumBuf.Get<float>();
	173	+ AscendC::LocalTensor<float> reduceWork;
	174	+ if constexpr (std::is_same_v<T, float>) {
	175	+ reduceWork = tmpBuffer.Get<float>();
	176	+ } else {
	177	+ reduceWork = tmpFloat.Get<float>();
	178	+ }
	179	+
	180	+ // 使用 Init 阶段预计算的 lastTileValidLen
	181	+ const uint32_t curLastTileValidLen = this->lastTileValidLen;
	182	+
	183	+ // --- ping-pong 双缓冲：循环前预取第 0 个 tile ---
	184	+ CopyIn(0);
	185	+
	186	+ // 前 loopCount-1 个完整 tile，无需判断是否为最后一个
	187	+ for (uint32_t t = 0; t + 1 < loopCount; ++t) {
	188	+ CopyIn(t + 1);
	189	+
	190	+ AscendC::LocalTensor<T> tileLocal = inQueueInput.DeQue<T>();
	191	+
	192	+ AscendC::LocalTensor<float> tileFloat;
	193	+ if constexpr (std::is_same_v<T, float>) {
	194	+ tileFloat = tileLocal.template ReinterpretCast<float>();
	195	+ } else {
	196	+ tileFloat = tmpFloat.Get<float>();
	197	+ AscendC::Cast(tileFloat, tileLocal, AscendC::RoundMode::CAST_NONE, tileLen);
	198	+ }
	199	+
	200	+ AscendC::Muls(tileFloat, tileFloat, INV_SQRT2, tileLen);
	201	+ AscendC::Mul(tileFloat, tileFloat, tileFloat, tileLen);
	202	+ AscendC::ReduceSum(tileSumTensor, tileFloat, reduceWork, tileLen);
	203	+ localSum += tileSumTensor.GetValue(0);
	204	+
	205	+ inQueueInput.FreeTensor(tileLocal);
	206	+ }
	207	+
	208	+ // 最后一个 tile 单独处理：长度为 lastTileLen，可能含 padding 需清零
	209	+ {
	210	+ AscendC::LocalTensor<T> tileLocal = inQueueInput.DeQue<T>();
	211	+
	212	+ AscendC::LocalTensor<float> tileFloat;
	213	+ if constexpr (std::is_same_v<T, float>) {
	214	+ tileFloat = tileLocal.template ReinterpretCast<float>();
	215	+ } else {
	216	+ tileFloat = tmpFloat.Get<float>();
	217	+ AscendC::Cast(tileFloat, tileLocal, AscendC::RoundMode::CAST_NONE, lastTileLen);
	218	+ }
	219	+
	220	+ // DataCopyPad 已在搬运时对 padding 区域补零，无需手动清零
	221	+ (void)curLastTileValidLen;
	222	+
	223	+ AscendC::Muls(tileFloat, tileFloat, INV_SQRT2, lastTileLen);
	224	+ AscendC::Mul(tileFloat, tileFloat, tileFloat, lastTileLen);
	225	+ AscendC::ReduceSum(tileSumTensor, tileFloat, reduceWork, lastTileLen);
	226	+ localSum += tileSumTensor.GetValue(0);
	227	+
	228	+ inQueueInput.FreeTensor(tileLocal);
	229	+ }
	230	+
	231	+ AscendC::LocalTensor<float> tmpBuf = tmpBuffer.Get<float>();
	232	+ AscendC::Duplicate(tmpBuf, 0.0f, SLOT_STRIDE);
	233	+ tmpBuf.SetValue(0, localSum);
	於於於欣洁6月4日 `tmpBuf.SetValue(0, localSum)` 只写入了第一个值，但 `DataCopy(workGm, tmpBuf, SLOT_STRIDE)` 会将 16 个 float 参与 atomic add。`tmpBuffer` 其余 15 个 float 未初始化，可能污染 workspace 同 cache line 内数据。
	234	+
	235	+ if constexpr (schMode == ELEMENTWISE_TPL_SCH_MODE_0) {
	236	+ if constexpr (std::is_same_v<T, float>) {
	237	+ AscendC::DataCopyExtParams outParams{1, static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
	238	+ AscendC::DataCopyPad(zGm, tmpBuf, outParams);
	239	+ } else {
	240	+ AscendC::LocalTensor<T> dstBuf = tileSumBuf.Get<T>();
	241	+ AscendC::Cast(dstBuf, tmpBuf, AscendC::RoundMode::CAST_ROUND, 16);
	242	+ AscendC::DataCopyExtParams outParams{1, static_cast<uint32_t>(sizeof(T)), 0, 0, 0};
	243	+ AscendC::DataCopyPad(zGm, dstBuf, outParams);
	244	+ }
	245	+ } else {
	246	+ AscendC::SetAtomicAdd<float>();
	247	+ AscendC::DataCopy(workGm, tmpBuf, SLOT_STRIDE);
	248	+ AscendC::SetAtomicNone();
	249	+ AscendC::DataCacheCleanAndInvalid<float, AscendC::CacheLine::SINGLE_CACHE_LINE, AscendC::DcciDst::CACHELINE_OUT>(workGm[0]);
	250	+ AscendC::SyncAll<true>();
	251	+
	252	+ if (this->blockIdx == 0) {
	253	+ float globalSum = workGm.GetValue(0);
	254	+ AscendC::LocalTensor<float> srcBuf = tmpBuffer.Get<float>();
	255	+ srcBuf.SetValue(0, globalSum);
	256	+ if constexpr (std::is_same_v<T, float>) {
	257	+ AscendC::DataCopyExtParams outParams{1, static_cast<uint32_t>(sizeof(float)), 0, 0, 0};
	258	+ AscendC::DataCopyPad(zGm, srcBuf, outParams);
	259	+ } else {
	260	+ AscendC::LocalTensor<T> dstBuf = tileSumBuf.Get<T>();
	261	+ AscendC::Cast(dstBuf, srcBuf, AscendC::RoundMode::CAST_ROUND, 16);
	262	+ AscendC::DataCopyExtParams outParams{1, static_cast<uint32_t>(sizeof(T)), 0, 0, 0};
	263	+ AscendC::DataCopyPad(zGm, dstBuf, outParams);
	264	+ }
	265	+ }
	266	+ }
	267	+}
	268	+} // namespace NsL2Loss
	269	+#endif // _L2_LOSS_H

Aexperimental/loss/l2_loss/op_kernel/l2_loss_tiling_data.h+30-0

@@ -0,0 +1,30 @@
	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 l2_loss_tiling_data.h
	12	+ * \brief tiling data struct
	13	+ */
	14	+
	15	+#ifndef L2_LOSS_TILING_DATA_H
	16	+#define L2_LOSS_TILING_DATA_H
	17	+
	18	+struct L2LossTilingData {
	19	+ uint32_t smallCoreDataNum;
	20	+ uint32_t bigCoreDataNum;
	21	+ uint32_t finalBigTileNum;
	22	+ uint32_t finalSmallTileNum;
	23	+ uint32_t tileDataNum;
	24	+ uint32_t smallTailDataNum;
	25	+ uint32_t bigTailDataNum;
	26	+ uint32_t tailBlockNum;
	27	+ uint32_t inputNum;
	28	+ uint32_t blockNum; // actual number of cores used for reduction in core 0
	29	+};
	30	+#endif // L2_LOSS_TILING_DATA_H

Aexperimental/loss/l2_loss/op_kernel/l2_loss_tiling_key.h+41-0

@@ -0,0 +1,41 @@
	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 l2_loss_tiling_key.h
	12	+ * \brief l2_loss tiling key declarations
	13	+*/
	14	+
	15	+#ifndef L2_LOSS_TILING_KEY_H
	16	+#define L2_LOSS_TILING_KEY_H
	17	+
	18	+#include "ascendc/host_api/tiling/template_argument.h"
	19	+
	20	+#define ELEMENTWISE_TPL_SCH_MODE_0 0
	21	+#define ELEMENTWISE_TPL_SCH_MODE_1 1
	22	+
	23	+ASCENDC_TPL_ARGS_DECL(L2Loss,
	24	+ ASCENDC_TPL_UINT_DECL(schMode, 1,
	25	+ ASCENDC_TPL_UI_LIST,
	26	+ ELEMENTWISE_TPL_SCH_MODE_0,
	27	+ ELEMENTWISE_TPL_SCH_MODE_1));
	28	+
	29	+ASCENDC_TPL_SEL(
	30	+ ASCENDC_TPL_ARGS_SEL(
	31	+ ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_AIV_ONLY),
	32	+ ASCENDC_TPL_UINT_SEL(schMode,
	33	+ ASCENDC_TPL_UI_LIST,
	34	+ ELEMENTWISE_TPL_SCH_MODE_0)),
	35	+ ASCENDC_TPL_ARGS_SEL(
	36	+ ASCENDC_TPL_KERNEL_TYPE_SEL(ASCENDC_TPL_MIX_AIV_1_0),
	37	+ ASCENDC_TPL_UINT_SEL(schMode,
	38	+ ASCENDC_TPL_UI_LIST,
	39	+ ELEMENTWISE_TPL_SCH_MODE_1)));
	40	+
	41	+#endif // L2_LOSS_TILING_KEY_H

Aexperimental/loss/l2_loss/tests/CMakeLists.txt+16-0

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

Aexperimental/loss/l2_loss/tests/ut/CMakeLists.txt+16-0

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

Aexperimental/loss/l2_loss/tests/ut/op_host/CMakeLists.txt+21-0

@@ -0,0 +1,21 @@
	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(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
	13	+ add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
	14	+endif()
	15	+
	16	+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
	17	+foreach(SUB_DIR ${CURRENT_DIRS})
	18	+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
	19	+ add_subdirectory(${SUB_DIR})
	20	+ endif()
	21	+endforeach()

Aexperimental/loss/l2_loss/tests/ut/op_host/l2_loss_tiling.h+25-0

@@ -0,0 +1,25 @@
	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 l2_loss_tiling.h
	13	+ * \brief
	14	+ */
	15	+
	16	+#ifndef L2_LOSS_TILING_H
	17	+#define L2_LOSS_TILING_H
	18	+
	19	+#include "register/tilingdata_base.h"
	20	+
	21	+namespace optiling {
	22	+struct L2LossCompileInfo {};
	23	+} // namespace optiling
	24	+
	25	+#endif // L2_LOSS_TILING_H

Aexperimental/loss/l2_loss/tests/ut/op_host/test_l2_loss_infershape.cpp+72-0

@@ -0,0 +1,72 @@
	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	+#include <iostream>
	12	+#include <gtest/gtest.h>
	13	+#include "exe_graph/runtime/storage_format.h"
	14	+#include "exe_graph/runtime/storage_shape.h"
	15	+#include "infer_shape_context_faker.h"
	16	+#include "op_impl_registry.h"
	17	+
	18	+using namespace std;
	19	+using namespace ge;
	20	+
	21	+class L2LossInfershape : public testing::Test {
	22	+protected:
	23	+ static void SetUpTestCase() { cout << "L2LossInfershape SetUp" << endl; }
	24	+ static void TearDownTestCase() { cout << "L2LossInfershape TearDown" << endl; }
	25	+};
	26	+
	27	+static ge::graphStatus RunL2LossInfershape(const gert::StorageShape& xShape, ge::DataType dtype,
	28	+ int64_t& outDim0)
	29	+{
	30	+ auto infershape_func = gert::OpImplRegistry::GetInstance().GetOpImpl("L2Loss")->infer_shape;
	31	+ auto holder = gert::InferShapeContextFaker()
	32	+ .SetOpType("L2Loss")
	33	+ .NodeIoNum(1, 1)
	34	+ .NodeInputTd(0, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
	35	+ .NodeOutputTd(0, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
	36	+ .InputShapes({const_cast<gert::StorageShape*>(&xShape)})
	37	+ .Build();
	38	+ gert::InferShapeContext* context = holder.GetContext<gert::InferShapeContext>();
	39	+ if (context == nullptr) { return ge::GRAPH_FAILED; }
	40	+ ge::graphStatus ret = infershape_func(context);
	41	+ if (ret == ge::GRAPH_SUCCESS) {
	42	+ const gert::Shape* outShape = context->GetOutputShape(0);
	43	+ if (outShape != nullptr && outShape->GetDimNum() >= 1) {
	44	+ outDim0 = outShape->GetDim(0);
	45	+ }
	46	+ }
	47	+ return ret;
	48	+}
	49	+
	50	+TEST_F(L2LossInfershape, l2_loss_infershape_fp32_success)
	51	+{
	52	+ gert::StorageShape x_shape({3, 4}, {3, 4});
	53	+ int64_t outDim0 = -1;
	54	+ EXPECT_EQ(RunL2LossInfershape(x_shape, ge::DT_FLOAT, outDim0), ge::GRAPH_SUCCESS);
	55	+ EXPECT_EQ(outDim0, 1);
	56	+}
	57	+
	58	+TEST_F(L2LossInfershape, l2_loss_infershape_fp16_success)
	59	+{
	60	+ gert::StorageShape x_shape({5, 10}, {5, 10});
	61	+ int64_t outDim0 = -1;
	62	+ EXPECT_EQ(RunL2LossInfershape(x_shape, ge::DT_FLOAT16, outDim0), ge::GRAPH_SUCCESS);
	63	+ EXPECT_EQ(outDim0, 1);
	64	+}
	65	+
	66	+TEST_F(L2LossInfershape, l2_loss_infershape_dynamic_shape)
	67	+{
	68	+ gert::StorageShape x_shape({5, -1}, {5, -1});
	69	+ int64_t outDim0 = -1;
	70	+ EXPECT_EQ(RunL2LossInfershape(x_shape, ge::DT_FLOAT, outDim0), ge::GRAPH_SUCCESS);
	71	+ EXPECT_EQ(outDim0, 1);
	72	+}

Aexperimental/loss/l2_loss/tests/ut/op_host/test_l2_loss_tiling.cpp+116-0

@@ -0,0 +1,116 @@
	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	+#include <iostream>
	12	+#include <vector>
	13	+#include <gtest/gtest.h>
	14	+#include "log/log.h"
	15	+#include "kernel_run_context_facker.h"
	16	+#include "exe_graph/runtime/storage_format.h"
	17	+#include "exe_graph/runtime/storage_shape.h"
	18	+#include "test_cube_util.h"
	19	+#include "register/op_impl_registry.h"
	20	+#include "ut_op_util.h"
	21	+#include "ut_op_common.h"
	22	+#include "platform/platform_infos_def.h"
	23	+
	24	+using namespace ut_util;
	25	+using namespace std;
	26	+using namespace ge;
	27	+
	28	+class L2LossTiling : public testing::Test {
	29	+protected:
	30	+ static void SetUpTestCase()
	31	+ {
	32	+ cout << "L2LossTiling SetUp" << endl;
	33	+ }
	34	+
	35	+ static void TearDownTestCase()
	36	+ {
	37	+ cout << "L2LossTiling TearDown" << endl;
	38	+ }
	39	+};
	40	+
	41	+static const string kCompileInfoString = R"({
	42	+ "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
	43	+ "Intrinsic_fix_pipe_l0c2out": false, "Intrinsic_data_move_l12ub": true, "Intrinsic_data_move_l0c2ub": true, "Intrinsic_data_move_out2l1_nd2nz": false,
	44	+ "UB_SIZE": 196608, "L2_SIZE": 33554432, "L1_SIZE": 524288,
	45	+ "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
	46	+ "CORE_NUM": 48}
	47	+ })";
	48	+
	49	+struct L2LossTilingCompileInfo {};
	50	+
	51	+static ge::graphStatus RunL2LossTiling(ge::DataType dtype)
	52	+{
	53	+ gert::StorageShape x_shape = {{1, 64, 2, 64}, {1, 64, 2, 64}};
	54	+ gert::StorageShape y_shape = {{1}, {1}};
	55	+
	56	+ map<string, string> soc_infos;
	57	+ map<string, string> aicore_spec;
	58	+ map<string, string> intrinsics;
	59	+ GetPlatFormInfos(kCompileInfoString.c_str(), soc_infos, aicore_spec, intrinsics);
	60	+
	61	+ fe::PlatFormInfos platform_info;
	62	+ platform_info.Init();
	63	+
	64	+ L2LossTilingCompileInfo compile_info;
	65	+
	66	+ auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl("L2Loss")->tiling;
	67	+
	68	+ auto kernel_holder =
	69	+ gert::KernelRunContextFaker()
	70	+ .KernelIONum(1, 1)
	71	+ .Inputs({const_cast<char>(kCompileInfoString.c_str()), reinterpret_cast<void>(&platform_info)})
	72	+ .Outputs({&compile_info})
	73	+ .Build();
	74	+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
	75	+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
	76	+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
	77	+ kernel_holder.GetContext<gert::TilingParseContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap",
	78	+ intrinsics);
	79	+
	80	+ auto param = gert::TilingData::CreateCap(4096);
	81	+ auto workspace_size_holder = gert::ContinuousVector::Create<size_t>(4096);
	82	+ auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holder.get());
	83	+ auto holder = gert::TilingContextFaker()
	84	+ .SetOpType("L2Loss")
	85	+ .NodeIoNum(1, 1)
	86	+ .IrInstanceNum({1})
	87	+ .InputShapes({&x_shape})
	88	+ .OutputShapes({&y_shape})
	89	+ .CompileInfo(&compile_info)
	90	+ .PlatformInfo(reinterpret_cast<char*>(&platform_info))
	91	+ .NodeInputTd(0, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
	92	+ .NodeOutputTd(0, dtype, ge::FORMAT_ND, ge::FORMAT_ND)
	93	+ .TilingData(param.get())
	94	+ .Workspace(ws_size)
	95	+ .Build();
	96	+ gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
	97	+ if (tiling_context == nullptr) {
	98	+ return ge::GRAPH_FAILED;
	99	+ }
	100	+ return tiling_func(tiling_context);
	101	+}
	102	+
	103	+TEST_F(L2LossTiling, l2_loss_tiling_fp32_success)
	104	+{
	105	+ EXPECT_EQ(RunL2LossTiling(ge::DT_FLOAT), ge::GRAPH_SUCCESS);
	106	+}
	107	+
	108	+TEST_F(L2LossTiling, l2_loss_tiling_fp16_success)
	109	+{
	110	+ EXPECT_EQ(RunL2LossTiling(ge::DT_FLOAT16), ge::GRAPH_SUCCESS);
	111	+}
	112	+
	113	+TEST_F(L2LossTiling, l2_loss_tiling_bf16_success)
	114	+{
	115	+ EXPECT_EQ(RunL2LossTiling(ge::DT_BF16), ge::GRAPH_SUCCESS);
	116	+}

Aexperimental/loss/l2_loss/tests/ut/op_kernel/CMakeLists.txt+23-0

@@ -0,0 +1,23 @@
	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_KERNEL_UT) AND NOT UT_DONE)
	12	+ # 算子自己的tiling文件路径
	13	+ set(l2_loss_tiling_files
	14	+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/l2_loss_tiling.cpp
	15	+ ${CMAKE_CURRENT_SOURCE_DIR}/../../../op_host/l2_loss_infershape.cpp
	16	+ )
	17	+ # 使用AddOpTestCase
	18	+ # param1：算子名称，以kernel方式命名
	19	+ # param2：soc版本，多个以分号分隔，例如："ascend910_95;ascend910b"
	20	+ # param3：自定义编译选项，一般填写测试的一种典型数据类型组合，不需要则传入空字符串
	21	+ # param4：该算子依赖的所有tiling源码文件
	22	+ AddOpTestCase(l2_loss "ascend910B1" "-DDTYPE_X=float" "${l2_loss_tiling_files}")
	23	+endif()

Aexperimental/loss/l2_loss/tests/ut/op_kernel/l2_loss_data/compare_data.py+59-0

@@ -0,0 +1,59 @@
	1	+#!/usr/bin/env python3
	2	+# -- coding: utf-8 --
	3	+# ----------------------------------------------------------------------------
	4	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	5	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	6	+# CANN Open Software License Agreement Version 2.0 (the "License").
	7	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	8	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	9	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	10	+# See LICENSE in the root of the software repository for the full text of the License.
	11	+# ----------------------------------------------------------------------------
	12	+
	13	+import sys
	14	+import numpy as np
	15	+import glob
	16	+import os
	17	+
	18	+curr_dir = os.path.dirname(os.path.realpath(__file__))
	19	+
	20	+def compare_data(golden_file_lists, output_file_lists, d_type):
	21	+ if d_type == "float16":
	22	+ np_dtype = np.float16
	23	+ elif d_type == "float32":
	24	+ np_dtype = np.float32
	25	+ else:
	26	+ raise ValueError("d_type must be float16 or float32")
	27	+
	28	+ data_same = True
	29	+ for gold, out in zip(golden_file_lists, output_file_lists):
	30	+ tmp_out = np.fromfile(out, np_dtype)
	31	+ tmp_gold = np.fromfile(gold, np_dtype)
	32	+ diff_res = np.isclose(tmp_out, tmp_gold, rtol=1e-3, atol=1e-5, equal_nan=True)
	33	+ diff_idx = np.where(diff_res != True)[0]
	34	+ if len(diff_idx) == 0:
	35	+ print("PASSED!")
	36	+ else:
	37	+ print("FAILED!")
	38	+ for idx in diff_idx[:5]:
	39	+ abs_diff = abs(float(tmp_out[idx]) - float(tmp_gold[idx]))
	40	+ rel_diff = abs_diff / (abs(float(tmp_gold[idx])) + 1e-12)
	41	+ print(f"index: [{idx}] real output value: [{tmp_out[idx]}] expected value: [{tmp_gold[idx]}]"
	42	+ f" abs_diff={abs_diff:.6g} rel_diff={rel_diff:.3e}")
	43	+ data_same = False
	44	+ return data_same
	45	+
	46	+def get_file_lists(dtype):
	47	+ golden_file_lists = sorted(glob.glob(curr_dir + "/golden.bin"))
	48	+ output_file_lists = sorted(glob.glob(curr_dir + "/output.bin"))
	49	+ return golden_file_lists, output_file_lists
	50	+
	51	+def process(d_type):
	52	+ golden_file_lists, output_file_lists = get_file_lists(d_type)
	53	+ result = compare_data(golden_file_lists, output_file_lists, d_type)
	54	+ print("compare result:", result)
	55	+ return result
	56	+
	57	+if __name__ == '__main__':
	58	+ ret = process(sys.argv[1])
	59	+ exit(0 if ret else 1)

Aexperimental/loss/l2_loss/tests/ut/op_kernel/l2_loss_data/gen_data.py+55-0

@@ -0,0 +1,55 @@
	1	+#!/usr/bin/env python3
	2	+# -- coding: utf-8 --
	3	+# ----------------------------------------------------------------------------
	4	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	5	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	6	+# CANN Open Software License Agreement Version 2.0 (the "License").
	7	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	8	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	9	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	10	+# See LICENSE in the root of the software repository for the full text of the License.
	11	+# ----------------------------------------------------------------------------
	12	+
	13	+import sys
	14	+import os
	15	+import numpy as np
	16	+import tensorflow as tf
	17	+
	18	+# 与算子内核保持完全一致的计算路径：先乘 1/√2 缩放，再平方，再求和
	19	+# 这是规定做法，golden 必须与算子路径对齐，否则因不同计算顺序导致的 ULP 差异会被当成算子错误
	20	+INV_SQRT2 = np.float32(0.70710678118654752)
	21	+
	22	+
	23	+def parse_str_to_shape_list(shape_str):
	24	+ shape_str = shape_str.strip('(').strip(')')
	25	+ shape_list = [int(x) for x in shape_str.split(",")]
	26	+ return np.array(shape_list)
	27	+
	28	+
	29	+def gen_data_and_golden(shape_str, d_type="float32"):
	30	+ d_type_dict = {
	31	+ "float32": np.float32,
	32	+ "float16": np.float16,
	33	+ "bfloat16": tf.bfloat16.as_numpy_dtype
	34	+ }
	35	+ np_type = d_type_dict[d_type]
	36	+ shape = parse_str_to_shape_list(shape_str)
	37	+ size = np.prod(shape)
	38	+ tmp_input = np.random.choice([0, 0.5, 1, 2, 3, 10, 100], size=size)
	39	+ tmp_input = tmp_input.reshape(shape).astype(np_type)
	40	+ # L2Loss golden：与算子内核完全一致的路径 —— (x * INV_SQRT2)^2 逐元素计算后求和
	41	+ # 用 float32 精度模拟算子行为，避免 golden 与算子因精度路径不同产生系统性偏差
	42	+ scaled = tmp_input.astype(np.float32) * INV_SQRT2
	43	+ tmp_golden = np.sum(scaled * scaled)
	44	+ tmp_golden = np.array([tmp_golden], dtype=np_type)
	45	+
	46	+ tmp_input.astype(np_type).tofile(f"{d_type}_input_t_l2_loss.bin")
	47	+ tmp_golden.astype(np_type).tofile(f"{d_type}_golden_t_l2_loss.bin")
	48	+
	49	+
	50	+if __name__ == "__main__":
	51	+ if len(sys.argv) != 3:
	52	+ print("Param num must be 3.")
	53	+ exit(1)
	54	+ os.system("rm -rf *.bin")
	55	+ gen_data_and_golden(sys.argv[1], sys.argv[2])

Aexperimental/loss/l2_loss/tests/ut/op_kernel/test_l2_loss.cpp+195-0

@@ -0,0 +1,195 @@
	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 test_l2_loss.cpp
	13	+ * \brief
	14	+ */
	15	+
	16	+#include <array>
	17	+#include <vector>
	18	+#include <iostream>
	19	+#include <string>
	20	+#include <cstdint>
	21	+#include "gtest/gtest.h"
	22	+
	23	+#ifdef __CCE_KT_TEST__
	24	+#include "tikicpulib.h"
	25	+#include "data_utils.h"
	26	+#endif
	27	+
	28	+#include "../../../op_kernel/l2_loss.cpp"
	29	+
	30	+using namespace std;
	31	+
	32	+extern "C" __global__ __aicore__ void l2_loss(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
	33	+
	34	+class L2LossTest : public testing::Test {
	35	+protected:
	36	+ static void SetUpTestCase()
	37	+ {
	38	+ std::cout << "l2_loss_test SetUp" << std::endl;
	39	+ const string cmd = "cp -rf " + dataPath + " ./";
	40	+ system(cmd.c_str());
	41	+ system("chmod -R 755 ./l2_loss_data/");
	42	+ }
	43	+ static void TearDownTestCase()
	44	+ {
	45	+ std::cout << "l2_loss_test TearDown" << std::endl;
	46	+ }
	47	+
	48	+private:
	49	+ const static std::string rootPath;
	50	+ const static std::string dataPath;
	51	+};
	52	+
	53	+const std::string L2LossTest::rootPath = "../../../../";
	54	+const std::string L2LossTest::dataPath = rootPath + "experimental/loss/l2_loss/tests/ut/op_kernel/l2_loss_data";
	55	+
	56	+template <typename T1, typename T2>
	57	+inline T1 CeilAlign(T1 a, T2 b)
	58	+{
	59	+ return (a + b - 1) / b * b;
	60	+}
	61	+
	62	+TEST_F(L2LossTest, test_case_float16_1)
	63	+{
	64	+ uint32_t blockDim = 1;
	65	+ uint32_t dataCount = 128 * 64;
	66	+ system("cd ./l2_loss_data/ && python3 gen_data.py '(128, 64)' 'float16'");
	67	+
	68	+ size_t inputByteSize = dataCount * sizeof(half);
	69	+ size_t outputByteSize = 1 * sizeof(half);
	70	+ std::string inputFileName = "./l2_loss_data/float16_input_t_l2_loss.bin";
	71	+
	72	+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));
	73	+ ReadFile(inputFileName, inputByteSize, x, inputByteSize);
	74	+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(CeilAlign(outputByteSize, 32));
	75	+
	76	+ size_t workspaceSize = 32 * 1024 * 1024;
	77	+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceSize);
	78	+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(L2LossTilingData));
	79	+
	80	+ L2LossTilingData* tilingData = reinterpret_cast<L2LossTilingData*>(tiling);
	81	+ tilingData->smallCoreDataNum = dataCount;
	82	+ tilingData->bigCoreDataNum = dataCount;
	83	+ tilingData->finalBigTileNum = 1;
	84	+ tilingData->finalSmallTileNum = 1;
	85	+ tilingData->tileDataNum = dataCount;
	86	+ tilingData->smallTailDataNum = dataCount;
	87	+ tilingData->bigTailDataNum = dataCount;
	88	+ tilingData->tailBlockNum = 0;
	89	+ tilingData->inputNum = dataCount;
	90	+ tilingData->blockNum = blockDim;
	91	+
	92	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	93	+ auto func = l2_loss<ELEMENTWISE_TPL_SCH_MODE_0>;
	94	+ ICPU_RUN_KF(func, blockDim, x, y, workspace, (uint8_t*)(tilingData));
	95	+
	96	+ std::string outputFileName = "./l2_loss_data/float16_output_t_l2_loss.bin";
	97	+ WriteFile(outputFileName, y, outputByteSize);
	98	+
	99	+ AscendC::GmFree((void*)(x));
	100	+ AscendC::GmFree((void*)(y));
	101	+ AscendC::GmFree((void*)workspace);
	102	+ AscendC::GmFree((void*)tiling);
	103	+
	104	+ system("cd ./l2_loss_data/ && python3 compare_data.py 'float16'");
	105	+}
	106	+
	107	+TEST_F(L2LossTest, test_case_float32_1)
	108	+{
	109	+ uint32_t blockDim = 1;
	110	+ uint32_t dataCount = 256 * 33;
	111	+ system("cd ./l2_loss_data/ && python3 gen_data.py '(256, 33)' 'float32'");
	112	+
	113	+ size_t inputByteSize = dataCount * sizeof(float);
	114	+ size_t outputByteSize = 1 * sizeof(float);
	115	+ std::string inputFileName = "./l2_loss_data/float32_input_t_l2_loss.bin";
	116	+
	117	+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));
	118	+ ReadFile(inputFileName, inputByteSize, x, inputByteSize);
	119	+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(CeilAlign(outputByteSize, 32));
	120	+
	121	+ size_t workspaceSize = 32 * 1024 * 1024;
	122	+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceSize);
	123	+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(L2LossTilingData));
	124	+
	125	+ L2LossTilingData* tilingData = reinterpret_cast<L2LossTilingData*>(tiling);
	126	+ tilingData->smallCoreDataNum = dataCount;
	127	+ tilingData->bigCoreDataNum = dataCount;
	128	+ tilingData->finalBigTileNum = 1;
	129	+ tilingData->finalSmallTileNum = 1;
	130	+ tilingData->tileDataNum = dataCount;
	131	+ tilingData->smallTailDataNum = dataCount;
	132	+ tilingData->bigTailDataNum = dataCount;
	133	+ tilingData->tailBlockNum = 0;
	134	+ tilingData->inputNum = dataCount;
	135	+ tilingData->blockNum = blockDim;
	136	+
	137	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	138	+ auto func = l2_loss<ELEMENTWISE_TPL_SCH_MODE_0>;
	139	+ ICPU_RUN_KF(func, blockDim, x, y, workspace, (uint8_t*)(tilingData));
	140	+
	141	+ std::string outputFileName = "./l2_loss_data/float32_output_t_l2_loss.bin";
	142	+ WriteFile(outputFileName, y, outputByteSize);
	143	+
	144	+ AscendC::GmFree((void*)(x));
	145	+ AscendC::GmFree((void*)(y));
	146	+ AscendC::GmFree((void*)workspace);
	147	+ AscendC::GmFree((void*)tiling);
	148	+
	149	+ system("cd ./l2_loss_data/ && python3 compare_data.py 'float32'");
	150	+}
	151	+
	152	+TEST_F(L2LossTest, test_case_bfloat16_1)
	153	+{
	154	+ uint32_t blockDim = 1;
	155	+ uint32_t dataCount = 128 * 64;
	156	+ system("cd ./l2_loss_data/ && python3 gen_data.py '(128, 64)' 'bfloat16'");
	157	+
	158	+ size_t inputByteSize = dataCount * sizeof(uint16_t); // bfloat16 is 2 bytes
	159	+ size_t outputByteSize = 1 * sizeof(uint16_t);
	160	+ std::string inputFileName = "./l2_loss_data/bfloat16_input_t_l2_loss.bin";
	161	+
	162	+ uint8_t* x = (uint8_t*)AscendC::GmAlloc(CeilAlign(inputByteSize, 32));
	163	+ ReadFile(inputFileName, inputByteSize, x, inputByteSize);
	164	+ uint8_t* y = (uint8_t*)AscendC::GmAlloc(CeilAlign(outputByteSize, 32));
	165	+
	166	+ size_t workspaceSize = 32 * 1024 * 1024;
	167	+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceSize);
	168	+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(L2LossTilingData));
	169	+
	170	+ L2LossTilingData* tilingData = reinterpret_cast<L2LossTilingData*>(tiling);
	171	+ tilingData->smallCoreDataNum = dataCount;
	172	+ tilingData->bigCoreDataNum = dataCount;
	173	+ tilingData->finalBigTileNum = 1;
	174	+ tilingData->finalSmallTileNum = 1;
	175	+ tilingData->tileDataNum = dataCount;
	176	+ tilingData->smallTailDataNum = dataCount;
	177	+ tilingData->bigTailDataNum = dataCount;
	178	+ tilingData->tailBlockNum = 0;
	179	+ tilingData->inputNum = dataCount;
	180	+ tilingData->blockNum = blockDim;
	181	+
	182	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	183	+ auto func = l2_loss<ELEMENTWISE_TPL_SCH_MODE_0>;
	184	+ ICPU_RUN_KF(func, blockDim, x, y, workspace, (uint8_t*)(tilingData));
	185	+
	186	+ std::string outputFileName = "./l2_loss_data/bfloat16_output_t_l2_loss.bin";
	187	+ WriteFile(outputFileName, y, outputByteSize);
	188	+
	189	+ AscendC::GmFree((void*)(x));
	190	+ AscendC::GmFree((void*)(y));
	191	+ AscendC::GmFree((void*)workspace);
	192	+ AscendC::GmFree((void*)tiling);
	193	+
	194	+ system("cd ./l2_loss_data/ && python3 compare_data.py 'bfloat16'");
	195	+}