Star921
1841
代码
代码
Issues231
Pull Requests361
Actions
项目讨论
Wiki
分析
Star921
1841
已合并
AdaLayerNormV2、AdaLayerNormQuant支持A5和性能优化 #912
陈海杰创建于 1月22日
AdaLayerNormV2、AdaLayerNormQuant支持A5和性能优化 #912
已合并
陈海杰创建于 1月22日
从chenhaijie1423/ops-nnAdaLayerNorm合入到cann/ops-nnmaster
共 28 个文件变更+3353-439
Mnorm/ada_layer_norm/op_host/ada_layer_norm_tiling.cpp+57-9
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -15,9 +15,12 @@
15#include "register/op_def_registry.h"15#include "register/op_def_registry.h"
16#include "log/log.h"16#include "log/log.h"
17#include "error_util.h"17#include "error_util.h"
18+#include "tiling_base/tiling_util.h"
18#include "ada_layer_norm_tiling.h"19#include "ada_layer_norm_tiling.h"
19 20 
20namespace optiling {21namespace optiling {
22+using namespace Ops::NN::OpTiling;
23+ 
21constexpr int32_t INPUT_TENSOR_NUM = 3;24constexpr int32_t INPUT_TENSOR_NUM = 3;
22constexpr int32_t INDEX_ZERO = 0;25constexpr int32_t INDEX_ZERO = 0;
23constexpr int32_t INDEX_ONE = 1;26constexpr int32_t INDEX_ONE = 1;
@@ -25,10 +28,12 @@ constexpr int32_t INDEX_TWO = 2;
25constexpr int32_t INDEX_THREE = 3;28constexpr int32_t INDEX_THREE = 3;
26constexpr int32_t INDEX_FOUR = 4;29constexpr int32_t INDEX_FOUR = 4;
27constexpr int32_t INDEX_FIVE = 5;30constexpr int32_t INDEX_FIVE = 5;
28-constexpr int32_t BLOCK_NUM = 32;31+constexpr int32_t BLOCK_SIZE = 32;
29 32 
30constexpr int64_t MULTI_ROW_SIZE = 3072;33constexpr int64_t MULTI_ROW_SIZE = 3072;
31constexpr int64_t SINGLE_ROW_SIZE = 6144;34constexpr int64_t SINGLE_ROW_SIZE = 6144;
35+constexpr uint32_t FP32_BYTE = 4;
36+constexpr uint32_t FP16_BYTE = 2;
32constexpr uint64_t WORK_SPACE_SIZE = 16 * 1024 * 1024;37constexpr uint64_t WORK_SPACE_SIZE = 16 * 1024 * 1024;
33 38 
34constexpr uint8_t BASE_OP_CODE = 1;39constexpr uint8_t BASE_OP_CODE = 1;
@@ -36,6 +41,7 @@ constexpr uint8_t BASE_V2_OP_CODE = 12;
36constexpr uint8_t QUANT_OP_CODE = 2;41constexpr uint8_t QUANT_OP_CODE = 2;
37constexpr uint8_t TILING_KEY_ONE = 1;42constexpr uint8_t TILING_KEY_ONE = 1;
38constexpr uint8_t TILING_KEY_TWO = 2;43constexpr uint8_t TILING_KEY_TWO = 2;
44+constexpr uint8_t TILING_KEY_FACTOR = 10;
39 45 
40class AdaLayerNormTiling46class AdaLayerNormTiling
41{47{
@@ -47,6 +53,8 @@ public:
47private:53private:
48 int32_t SplitCore(int32_t coreNumPlatform);54 int32_t SplitCore(int32_t coreNumPlatform);
49 void FillTilingData();55 void FillTilingData();
56+ void DoLayerNormTiling();
57+ uint8_t GetTilingKey(bool isRegBase);
50 template <typename T1, typename T2>58 template <typename T1, typename T2>
51 inline auto CeilA2B(T1 a, T2 b) const -> T1;59 inline auto CeilA2B(T1 a, T2 b) const -> T1;
52 60 
@@ -115,12 +123,16 @@ ge::graphStatus AdaLayerNormTiling::RunBigKernelTiling()
115 }123 }
116 seqLen = xShape.GetDim(xDim - INDEX_TWO);124 seqLen = xShape.GetDim(xDim - INDEX_TWO);
117 hiddenDim = xShape.GetDim(xDim - INDEX_ONE);125 hiddenDim = xShape.GetDim(xDim - INDEX_ONE);
118- hiddenDimCeil = CeilA2B(hiddenDim, BLOCK_NUM) * BLOCK_NUM;126+ hiddenDimCeil = CeilA2B(hiddenDim, BLOCK_SIZE) * BLOCK_SIZE;
119 epsilon = *tilingContext->GetAttrs()->GetAttrPointer<float>(0);127 epsilon = *tilingContext->GetAttrs()->GetAttrPointer<float>(0);
120 128 
121 auto compileInfo = reinterpret_cast<const AdaLayerNormCompileInfo*>(tilingContext->GetCompileInfo());129 auto compileInfo = reinterpret_cast<const AdaLayerNormCompileInfo*>(tilingContext->GetCompileInfo());
122 int32_t coreNumPlatform = compileInfo->coreNum;130 int32_t coreNumPlatform = compileInfo->coreNum;
123 int32_t needCoreNum = SplitCore(coreNumPlatform);131 int32_t needCoreNum = SplitCore(coreNumPlatform);
132+ if (compileInfo->isRegBase) {
133+ DoLayerNormTiling();
134+ }
135+ 
124 size_t* workspaces = tilingContext->GetWorkspaceSizes(1);136 size_t* workspaces = tilingContext->GetWorkspaceSizes(1);
125 if (opCode == QUANT_OP_CODE && hiddenDim > SINGLE_ROW_SIZE) {137 if (opCode == QUANT_OP_CODE && hiddenDim > SINGLE_ROW_SIZE) {
126 workspaces[0] = WORK_SPACE_SIZE + needCoreNum * hiddenDimCeil * sizeof(float);138 workspaces[0] = WORK_SPACE_SIZE + needCoreNum * hiddenDimCeil * sizeof(float);
@@ -128,12 +140,8 @@ ge::graphStatus AdaLayerNormTiling::RunBigKernelTiling()
128 workspaces[0] = WORK_SPACE_SIZE;140 workspaces[0] = WORK_SPACE_SIZE;
129 }141 }
130 142 
131- tilingContext->SetBlockDim(coreNumPlatform);143+ tilingContext->SetBlockDim(needCoreNum);
132- if (opCode == BASE_V2_OP_CODE && isWeightFloat) {144+ tilingContext->SetTilingKey(GetTilingKey(compileInfo->isRegBase));
133- tilingContext->SetTilingKey(TILING_KEY_TWO);
134- } else {
135- tilingContext->SetTilingKey(TILING_KEY_ONE);
136- }
137 FillTilingData();145 FillTilingData();
138 return ge::GRAPH_SUCCESS;146 return ge::GRAPH_SUCCESS;
139}147}
@@ -147,6 +155,7 @@ int32_t AdaLayerNormTiling::SplitCore(int32_t coreNumPlatform)
147 } else if (hiddenDim > SINGLE_ROW_SIZE) {155 } else if (hiddenDim > SINGLE_ROW_SIZE) {
148 int64_t batch = CeilA2B(hiddenDim, SINGLE_ROW_SIZE);156 int64_t batch = CeilA2B(hiddenDim, SINGLE_ROW_SIZE);
149 sliceSize = CeilA2B(hiddenDim, batch);157 sliceSize = CeilA2B(hiddenDim, batch);
158+ sliceSize = CeilA2B(sliceSize, BLOCK_SIZE) * BLOCK_SIZE;
150 }159 }
151 160 
152 int64_t singleCoreNum = coreNumPlatform != 0 ? (batchSize * seqLen) / coreNumPlatform : 0;161 int64_t singleCoreNum = coreNumPlatform != 0 ? (batchSize * seqLen) / coreNumPlatform : 0;
@@ -158,6 +167,41 @@ int32_t AdaLayerNormTiling::SplitCore(int32_t coreNumPlatform)
158 return singleCoreNum > 0 ? coreNumPlatform : tailNum;167 return singleCoreNum > 0 ? coreNumPlatform : tailNum;
159}168}
160 169 
170+uint8_t AdaLayerNormTiling::GetTilingKey(bool isRegBase)
171+{
172+ uint8_t tilingKey = 0;
173+ if (isRegBase) {
174+ tilingKey = (hiddenDim > SINGLE_ROW_SIZE) ? TILING_KEY_TWO : TILING_KEY_ONE;
175+ tilingKey *= TILING_KEY_FACTOR;
176+ }
177+ tilingKey += (opCode == BASE_V2_OP_CODE && isWeightFloat) ? TILING_KEY_TWO : TILING_KEY_ONE;
178+ return tilingKey;
179+}
180+ 
181+void AdaLayerNormTiling::DoLayerNormTiling()
182+{
183+ uint32_t minValue;
184+ uint32_t maxValue;
185+ if (hiddenDim > SINGLE_ROW_SIZE) {
186+ int64_t sliceSize = tilingData.get_sliceSize();
187+ int64_t tmpBufferSize = 0;
188+ ge::Shape inputShape({1, sliceSize});
189+ uint32_t welfordXByte = (dataType == ge::DataType::DT_FLOAT) ? FP32_BYTE : FP16_BYTE;
190+ AscendC::GetWelfordUpdateMaxMinTmpSize(inputShape, welfordXByte, FP32_BYTE, false, true, maxValue, minValue);
191+ tmpBufferSize += maxValue;
192+ AscendC::GetWelfordFinalizeMaxMinTmpSize(inputShape, FP32_BYTE, false, maxValue, minValue);
193+ tmpBufferSize += maxValue;
194+ AscendC::GetNormalizeMaxMinTmpSize(inputShape, FP32_BYTE, FP32_BYTE, false, true, false, maxValue, minValue);
195+ tmpBufferSize += maxValue;
196+ tilingData.set_tmpBufferSize(tmpBufferSize);
197+ } else {
198+ ge::Shape inputShape({tilingData.get_rowNum(), hiddenDim});
199+ AscendC::GetLayerNormMaxMinTmpSize(inputShape, FP32_BYTE, true, true, false, maxValue, minValue);
200+ AscendC::GetLayerNormNDTilingInfo(inputShape, maxValue, FP32_BYTE, true, true, tilingData.layerNormTiling);
201+ tilingData.set_tmpBufferSize(maxValue);
202+ }
203+}
204+ 
161void AdaLayerNormTiling::FillTilingData()205void AdaLayerNormTiling::FillTilingData()
162{206{
163 tilingData.set_batchSize(batchSize);207 tilingData.set_batchSize(batchSize);
@@ -213,6 +257,10 @@ static ge::graphStatus TilingPrepareTiling(gert::TilingParseContext* context)
213 OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);257 OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
214 auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());258 auto ascendcPlatform = platform_ascendc::PlatformAscendC(context->GetPlatformInfo());
215 compileInfo->coreNum = ascendcPlatform.GetCoreNumAiv();259 compileInfo->coreNum = ascendcPlatform.GetCoreNumAiv();
260+ uint64_t ubSizePlatForm = 0;
261+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatForm);
262+ compileInfo->ubSizePlatForm = ubSizePlatForm;
263+ compileInfo->isRegBase = IsRegbaseSocVersion(context);
216 264 
217 OP_TILING_CHECK(265 OP_TILING_CHECK(
218 compileInfo->coreNum <= 0,266 compileInfo->coreNum <= 0,
Mnorm/ada_layer_norm/op_host/ada_layer_norm_tiling.h+8-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -21,7 +21,9 @@
21namespace optiling {21namespace optiling {
22 22 
23struct AdaLayerNormCompileInfo {23struct AdaLayerNormCompileInfo {
24- int32_t coreNum;24+ int32_t coreNum = 0;
25+ uint64_t ubSizePlatForm = 0;
26+ bool isRegBase = false;
25};27};
26 28 
27BEGIN_TILING_DATA_DEF(AdaLayerNormTilingData)29BEGIN_TILING_DATA_DEF(AdaLayerNormTilingData)
@@ -40,6 +42,10 @@ TILING_DATA_FIELD_DEF(int64_t, tailNum)
40// 分批参数，sliceSize < hiddenDim说明需要分批处理，rowNum > 1说明一次处理多行数据42// 分批参数，sliceSize < hiddenDim说明需要分批处理，rowNum > 1说明一次处理多行数据
41TILING_DATA_FIELD_DEF(int64_t, sliceSize)43TILING_DATA_FIELD_DEF(int64_t, sliceSize)
42TILING_DATA_FIELD_DEF(int64_t, rowNum)44TILING_DATA_FIELD_DEF(int64_t, rowNum)
45+// 临时buffer大小
46+TILING_DATA_FIELD_DEF(int64_t, tmpBufferSize)
47+// LayerNormTiling
48+TILING_DATA_FIELD_DEF_STRUCT(LayerNormSeparateTiling, layerNormTiling)
43END_TILING_DATA_DEF;49END_TILING_DATA_DEF;
44 50 
45REGISTER_TILING_DATA_CLASS(AdaLayerNorm, AdaLayerNormTilingData)51REGISTER_TILING_DATA_CLASS(AdaLayerNorm, AdaLayerNormTilingData)
Mnorm/ada_layer_norm/op_kernel/ada_layer_norm.cpp+2-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -22,6 +22,7 @@ extern "C" __global__ __aicore__ void ada_layer_norm(
22 GM_ADDR x, GM_ADDR scale, GM_ADDR shift, GM_ADDR weight, GM_ADDR bias, GM_ADDR out, GM_ADDR workspace,22 GM_ADDR x, GM_ADDR scale, GM_ADDR shift, GM_ADDR weight, GM_ADDR bias, GM_ADDR out, GM_ADDR workspace,
23 GM_ADDR tiling)23 GM_ADDR tiling)
24{24{
25+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
25 GET_TILING_DATA(tilingData, tiling);26 GET_TILING_DATA(tilingData, tiling);
26 27 
27#define INIT_AND_PROCESS \28#define INIT_AND_PROCESS \
Mnorm/ada_layer_norm/op_kernel/ada_layer_norm_base.h+7-66
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -15,47 +15,11 @@
15#ifndef ADA_LAYER_NORM_BASE_H15#ifndef ADA_LAYER_NORM_BASE_H
16#define ADA_LAYER_NORM_BASE_H16#define ADA_LAYER_NORM_BASE_H
17 17 
18-#include "kernel_operator.h"18+#include "ada_layer_norm_util.h"
19 19 
20namespace AdaLayerNormNS {20namespace AdaLayerNormNS {
21using namespace AscendC;21using namespace AscendC;
22 22 
23-constexpr uint8_t BASE_OP_CODE = 1;
24-constexpr uint8_t BASE_V2_OP_CODE = 12;
25-constexpr uint8_t QUANT_OP_CODE = 2;
26-constexpr int32_t MAX_X_SIZE = 16384;
27-constexpr int32_t TENSOR_NUM = 7;
28-constexpr int32_t DATA_COUNT = 6144;
29-constexpr int32_t BATCH_COUNT = 1024;
30-constexpr int32_t INT8_BLOCK_NUM = 32;
31-constexpr int32_t HALF_BLOCK_NUM = 16;
32-constexpr int32_t FLOAT_BLOCK_NUM = 8;
33-constexpr float MAX_INT8 = 127.0f;
34-constexpr float ONE_FLOAT = 1.0f;
35-constexpr float FACTOR_INT8 = 1.0f / 127.0f;
36- 
37-struct RowRange {
38- int64_t rowStart;
39- int64_t rowEnd;
40- int64_t actualRowNum;
41- int64_t batchStart;
42- int64_t batchEnd;
43- int64_t dataCount;
44-};
45- 
46-struct GmAddr {
47- const GM_ADDR x = nullptr;
48- const GM_ADDR scale = nullptr;
49- const GM_ADDR shift = nullptr;
50- const GM_ADDR weight = nullptr;
51- const GM_ADDR bias = nullptr;
52- const GM_ADDR smooth_scales = nullptr;
53- const GM_ADDR out = nullptr;
54- const GM_ADDR mean = nullptr;
55- const GM_ADDR rstd = nullptr;
56- const GM_ADDR quant_scale = nullptr;
57-};
58- 
59template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>23template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>
60class AdaLayerNormND {24class AdaLayerNormND {
61public:25public:
@@ -66,18 +30,6 @@ public:
66 __aicore__ inline void Process();30 __aicore__ inline void Process();
67 31 
68private:32private:
69- template <typename T1, typename T2>
70- __aicore__ inline T1 CeilA2B(T1 a, T2 b)
71- {
72- return (b != 0) ? (a + b - 1) / b : a;
73- };
74- 
75- template <typename T1, typename T2>
76- __aicore__ inline T1 Min(T1 a, T2 b)
77- {
78- return (a < b) ? a : b;
79- };
80- 
81 __aicore__ inline void InitTensor();33 __aicore__ inline void InitTensor();
82 __aicore__ inline void InitEventId();34 __aicore__ inline void InitEventId();
83 __aicore__ inline void ReleaseEventId();35 __aicore__ inline void ReleaseEventId();
@@ -97,19 +49,13 @@ private:
97 __aicore__ inline void Adaption(RowRange range);49 __aicore__ inline void Adaption(RowRange range);
98 __aicore__ inline void DynamicQuant(RowRange range, int64_t batchIdx);50 __aicore__ inline void DynamicQuant(RowRange range, int64_t batchIdx);
99 51 
100- __aicore__ inline void CopyInData(LocalTensor<float> inputFloat, GlobalTensor<X_DTYPE> inputGm, int64_t len);
101- __aicore__ inline void CopyInWeightBias(LocalTensor<float> inputFloat, GlobalTensor<WEIGHT_DTYPE> inputGm, int64_t len);
102 __aicore__ inline void CopyInOtherData(int64_t offset, int64_t len);52 __aicore__ inline void CopyInOtherData(int64_t offset, int64_t len);
103 __aicore__ inline void CopyInScaleShift(int64_t offset, uint16_t blockCount, int64_t len);53 __aicore__ inline void CopyInScaleShift(int64_t offset, uint16_t blockCount, int64_t len);
104 __aicore__ inline void CopyInSlice(int64_t offset, int64_t scaleOffset, int64_t h, int64_t len);54 __aicore__ inline void CopyInSlice(int64_t offset, int64_t scaleOffset, int64_t h, int64_t len);
105- __aicore__ inline void CopyInSliceX(int64_t offset, int64_t len);
106 __aicore__ inline void CopyInX(int64_t offset, uint16_t blockCount, int64_t len);55 __aicore__ inline void CopyInX(int64_t offset, uint16_t blockCount, int64_t len);
107 __aicore__ inline void BaseCopyOut(int64_t offset, uint16_t blockCount, int64_t len);56 __aicore__ inline void BaseCopyOut(int64_t offset, uint16_t blockCount, int64_t len);
108- __aicore__ inline void QuantCopyOut(int64_t offset, uint16_t blockCount, int64_t len);
109 __aicore__ inline void CopyScaleOut(int64_t offset, int64_t len);57 __aicore__ inline void CopyScaleOut(int64_t offset, int64_t len);
110 __aicore__ inline void CopyMeanRstdOut(int64_t offset, int64_t len);58 __aicore__ inline void CopyMeanRstdOut(int64_t offset, int64_t len);
111- __aicore__ inline void CopyNormOut(int64_t h, int64_t len);
112- __aicore__ inline void CopyInNorm(int64_t h, int64_t len);
113 59 
114private:60private:
115 TPipe pipe;61 TPipe pipe;
@@ -150,7 +96,6 @@ private:
150 96 
151 event_t eventIdVToMte2;97 event_t eventIdVToMte2;
152 event_t eventIdMte2ToV;98 event_t eventIdMte2ToV;
153- event_t eventIdVToMte3;
154 event_t eventIdMte3ToV;99 event_t eventIdMte3ToV;
155 event_t eventIdVToS;100 event_t eventIdVToS;
156 event_t eventIdMte3ToS;101 event_t eventIdMte3ToS;
@@ -302,8 +247,6 @@ template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>
302__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::InitEventId()247__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::InitEventId()
303{248{
304 eventIdVToMte2 = static_cast<event_t>(pipe.AllocEventID<HardEvent::V_MTE2>());249 eventIdVToMte2 = static_cast<event_t>(pipe.AllocEventID<HardEvent::V_MTE2>());
305- eventIdMte2ToV = static_cast<event_t>(pipe.AllocEventID<HardEvent::MTE2_V>());
306- eventIdVToMte3 = static_cast<event_t>(pipe.AllocEventID<HardEvent::V_MTE3>());
307 eventIdMte3ToV = static_cast<event_t>(pipe.AllocEventID<HardEvent::MTE3_V>());250 eventIdMte3ToV = static_cast<event_t>(pipe.AllocEventID<HardEvent::MTE3_V>());
308 eventIdVToS = static_cast<event_t>(pipe.AllocEventID<HardEvent::V_S>());251 eventIdVToS = static_cast<event_t>(pipe.AllocEventID<HardEvent::V_S>());
309 eventIdMte3ToS = static_cast<event_t>(pipe.AllocEventID<HardEvent::MTE3_S>());252 eventIdMte3ToS = static_cast<event_t>(pipe.AllocEventID<HardEvent::MTE3_S>());
@@ -324,20 +267,20 @@ __aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::SliceProc
324 if (hiddenDim > MAX_X_SIZE) {267 if (hiddenDim > MAX_X_SIZE) {
325 for (int64_t h = 0; h < hiddenDim; h += MAX_X_SIZE) {268 for (int64_t h = 0; h < hiddenDim; h += MAX_X_SIZE) {
326 int64_t dataCount = Min(MAX_X_SIZE, hiddenDim - h);269 int64_t dataCount = Min(MAX_X_SIZE, hiddenDim - h);
327- CopyInSliceX(offset + h, dataCount);270+ CopyInAndCast(xFloat, xGm[offset + h], dataCount, MAX_X_SIZE);
328 ComputeMean(dataCount, meanValue);271 ComputeMean(dataCount, meanValue);
329 }272 }
330 for (int64_t h = 0; h < hiddenDim; h += MAX_X_SIZE) {273 for (int64_t h = 0; h < hiddenDim; h += MAX_X_SIZE) {
331 int64_t dataCount = Min(MAX_X_SIZE, hiddenDim - h);274 int64_t dataCount = Min(MAX_X_SIZE, hiddenDim - h);
332- CopyInSliceX(offset + h, dataCount);275+ CopyInAndCast(xFloat, xGm[offset + h], dataCount, MAX_X_SIZE);
333 ComputeVar(dataCount, meanValue, varValue);276 ComputeVar(dataCount, meanValue, varValue);
334 }277 }
335 } else {278 } else {
336- CopyInSliceX(offset, hiddenDim);279+ CopyInAndCast(xFloat, xGm[offset], hiddenDim, MAX_X_SIZE);
337 ComputeMean(hiddenDim, meanValue);280 ComputeMean(hiddenDim, meanValue);
338 ComputeVar(hiddenDim, meanValue, varValue);281 ComputeVar(hiddenDim, meanValue, varValue);
339 }282 }
340- float rstdValue = 1.0f / sqrt(varValue + epsilon);283+ float rstdValue = ONE_FLOAT / sqrt(varValue + epsilon);
341 284 
342 if constexpr (OP_CODE == QUANT_OP_CODE) {285 if constexpr (OP_CODE == QUANT_OP_CODE) {
343 QuantSliceCompute(rowIdx, batchCount, meanValue, rstdValue);286 QuantSliceCompute(rowIdx, batchCount, meanValue, rstdValue);
@@ -392,7 +335,7 @@ __aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::FastProce
392 Adaption(range);335 Adaption(range);
393 if constexpr (OP_CODE == QUANT_OP_CODE) {336 if constexpr (OP_CODE == QUANT_OP_CODE) {
394 DynamicQuant(range, batchCount);337 DynamicQuant(range, batchCount);
395- QuantCopyOut(range.rowStart * hiddenDim, range.actualRowNum, hiddenDim);338+ CopyOut(quantOutGm[range.rowStart * hiddenDim], yInt, range.actualRowNum, hiddenDim);
CANN-robot
CANN-robotCANN-robot1月22日

代码结构与可维护性: 第338行调用 CopyOut(quantOutGm[range.rowStart * hiddenDim], yInt, range.actualRowNum, hiddenDim);，但 CopyOut 函数未在类中声明或定义（仅在第56-59行有 BaseCopyOut、CopyScaleOut、CopyMeanRstdOut）。这可能是来自 ada_layer_norm_util.h 的通用函数，但在此上下文中使用不一致，降低了代码可读性和维护性。

问题类型: 代码结构与可维护性 文件路径: norm/ada_layer_norm/op_kernel/ada_layer_norm_base.h 行号: 338 问题代码:

            CopyOut(quantOutGm[range.rowStart * hiddenDim], yInt, range.actualRowNum, hiddenDim);

修改建议:

建议统一拷贝函数命名或添加类内声明。如果 `CopyOut` 是外部函数，应在头文件中添加注释说明其来源和用途，或考虑封装为类方法以保持接口一致性。

此评论由代码审查工具自动生成

likedislike
396 } else {339 } else {
397 BaseCopyOut(range.rowStart * hiddenDim, range.actualRowNum, hiddenDim);340 BaseCopyOut(range.rowStart * hiddenDim, range.actualRowNum, hiddenDim);
398 }341 }
@@ -417,8 +360,6 @@ template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>
417__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::ReleaseEventId()360__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::ReleaseEventId()
418{361{
419 pipe.ReleaseEventID<HardEvent::V_MTE2>(eventIdVToMte2);362 pipe.ReleaseEventID<HardEvent::V_MTE2>(eventIdVToMte2);
420- pipe.ReleaseEventID<HardEvent::MTE2_V>(eventIdMte2ToV);
421- pipe.ReleaseEventID<HardEvent::V_MTE3>(eventIdVToMte3);
422 pipe.ReleaseEventID<HardEvent::MTE3_V>(eventIdMte3ToV);363 pipe.ReleaseEventID<HardEvent::MTE3_V>(eventIdMte3ToV);
423 pipe.ReleaseEventID<HardEvent::V_S>(eventIdVToS);364 pipe.ReleaseEventID<HardEvent::V_S>(eventIdVToS);
424 pipe.ReleaseEventID<HardEvent::MTE3_S>(eventIdMte3ToS);365 pipe.ReleaseEventID<HardEvent::MTE3_S>(eventIdMte3ToS);
Mnorm/ada_layer_norm/op_kernel/ada_layer_norm_base_v1.h+19-110
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -90,22 +90,21 @@ __aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::QuantSlic
90 }90 }
91 Abs(yFloat, xFloat, dataCount);91 Abs(yFloat, xFloat, dataCount);
92 PipeBarrier<PIPE_V>();92 PipeBarrier<PIPE_V>();
93- CopyNormOut(h, dataCount);93+ CopyOut(normGm[normOffset + h], xFloat, 1, dataCount);
94 SetFlag<HardEvent::MTE3_MTE2>(eventIdMte3ToMte2);94 SetFlag<HardEvent::MTE3_MTE2>(eventIdMte3ToMte2);
95 WaitFlag<HardEvent::MTE3_MTE2>(eventIdMte3ToMte2);95 WaitFlag<HardEvent::MTE3_MTE2>(eventIdMte3ToMte2);
96- ReduceMax(reduceFloat, yFloat, yFloat, dataCount);96+ ReduceMaxCustom(reduceFloat, yFloat, dataCount);
97 SetFlag<HardEvent::V_S>(eventIdVToS);97 SetFlag<HardEvent::V_S>(eventIdVToS);
98 WaitFlag<HardEvent::V_S>(eventIdVToS);98 WaitFlag<HardEvent::V_S>(eventIdVToS);
99- float tmpMax = reduceFloat.GetValue(0);99+ maxValue = AscendC::Std::max(reduceFloat.GetValue(0), maxValue);
100- maxValue = (tmpMax != tmpMax) ? tmpMax : (tmpMax > maxValue ? tmpMax : maxValue);
101 }100 }
102 quantScaleFloat.SetValue(batchIdx, maxValue / MAX_INT8);101 quantScaleFloat.SetValue(batchIdx, maxValue / MAX_INT8);
103 // 计算量化输出102 // 计算量化输出
104 for (int64_t h = 0; h < hiddenDim; h += sliceSize) {103 for (int64_t h = 0; h < hiddenDim; h += sliceSize) {
105 int64_t dataCount = Min(sliceSize, hiddenDim - h);104 int64_t dataCount = Min(sliceSize, hiddenDim - h);
106- CopyInNorm(h, dataCount);105+ CopyInAndCast(xFloat, normGm[normOffset + h], dataCount, DATA_COUNT);
107 ComputeSliceQuant(maxValue, dataCount);106 ComputeSliceQuant(maxValue, dataCount);
108- QuantCopyOut(offset + h, 1, dataCount);107+ CopyOut(quantOutGm[offset + h], yInt, 1, dataCount);
109 }108 }
110}109}
111 110 
@@ -238,7 +237,7 @@ __aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::SingleLay
238 }237 }
239 SetFlag<HardEvent::V_S>(eventIdVToS);238 SetFlag<HardEvent::V_S>(eventIdVToS);
240 WaitFlag<HardEvent::V_S>(eventIdVToS);239 WaitFlag<HardEvent::V_S>(eventIdVToS);
241- float rstdValue = 1.0f / sqrt(reduceFloat.GetValue(0) + epsilon);240+ float rstdValue = ONE_FLOAT / sqrt(reduceFloat.GetValue(0) + epsilon);
242 if constexpr (OP_CODE == BASE_V2_OP_CODE) {241 if constexpr (OP_CODE == BASE_V2_OP_CODE) {
243 rstdOutFloat.SetValue(batchIdx, rstdValue);242 rstdOutFloat.SetValue(batchIdx, rstdValue);
244 }243 }
@@ -275,9 +274,9 @@ __aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::Adaption(
275 yOffset += hiddenDimCeil;274 yOffset += hiddenDimCeil;
276 }275 }
277 } else {276 } else {
278- Mul(yFloat, yFloat, scaleFloat, range.dataCount);277+ Mul(yFloat, yFloat, scaleFloat, hiddenDim);
279 PipeBarrier<PIPE_V>();278 PipeBarrier<PIPE_V>();
280- Add(yFloat, yFloat, shiftFloat, range.dataCount);279+ Add(yFloat, yFloat, shiftFloat, hiddenDim);
281 }280 }
282}281}
283 282 
@@ -295,8 +294,7 @@ __aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::DynamicQu
295 PipeBarrier<PIPE_V>();294 PipeBarrier<PIPE_V>();
296 if (range.actualRowNum > 1) {295 if (range.actualRowNum > 1) {
297 for (int64_t rowIdx = 0; rowIdx < range.actualRowNum; rowIdx++) {296 for (int64_t rowIdx = 0; rowIdx < range.actualRowNum; rowIdx++) {
298- ReduceMax(297+ ReduceMaxCustom(reduceFloat[rowIdx], xFloat[rowIdx * hiddenDimCeil], hiddenDim);
299- reduceFloat[rowIdx], xFloat[rowIdx * hiddenDimCeil], xFloat[rowIdx * hiddenDimCeil], hiddenDim, false);
300 }298 }
301 PipeBarrier<PIPE_V>();299 PipeBarrier<PIPE_V>();
302 SetFlag<HardEvent::V_S>(eventIdVToS);300 SetFlag<HardEvent::V_S>(eventIdVToS);
@@ -315,7 +313,7 @@ __aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::DynamicQu
315 SetFlag<HardEvent::V_MTE2>(eventIdVToMte2);313 SetFlag<HardEvent::V_MTE2>(eventIdVToMte2);
316 WaitFlag<HardEvent::V_MTE2>(eventIdVToMte2);314 WaitFlag<HardEvent::V_MTE2>(eventIdVToMte2);
317 } else {315 } else {
318- ReduceMax(reduceFloat, xFloat, xFloat, hiddenDim, false);316+ ReduceMaxCustom(reduceFloat, xFloat, hiddenDim);
319 SetFlag<HardEvent::V_MTE2>(eventIdVToMte2);317 SetFlag<HardEvent::V_MTE2>(eventIdVToMte2);
320 WaitFlag<HardEvent::V_MTE2>(eventIdVToMte2);318 WaitFlag<HardEvent::V_MTE2>(eventIdVToMte2);
321 SetFlag<HardEvent::V_S>(eventIdVToS);319 SetFlag<HardEvent::V_S>(eventIdVToS);
@@ -332,56 +330,18 @@ __aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::DynamicQu
332 }330 }
333}331}
334 332 
335-template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>
336-__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyInData(
337- LocalTensor<float> inputFloat, GlobalTensor<X_DTYPE> inputGm, int64_t len)
338-{
339- DataCopyExtParams copyInParams{1, static_cast<uint32_t>(len * sizeof(X_DTYPE)), 0, 0, 0};
340- DataCopyPadExtParams<X_DTYPE> padParams{false, 0, 0, 0};
341- if constexpr (std::is_same_v<X_DTYPE, float>) {
342- DataCopyPad(inputFloat, inputGm, copyInParams, padParams);
343- } else {
344- DataCopyPad(inputFloat.ReinterpretCast<X_DTYPE>()[DATA_COUNT], inputGm, copyInParams, padParams);
345- }
346- SetFlag<HardEvent::MTE2_V>(eventIdMte2ToV);
347- WaitFlag<HardEvent::MTE2_V>(eventIdMte2ToV);
348- if constexpr (!std::is_same_v<X_DTYPE, float>) {
349- Cast(inputFloat, inputFloat.ReinterpretCast<X_DTYPE>()[DATA_COUNT], RoundMode::CAST_NONE, len);
350- PipeBarrier<PIPE_V>();
351- }
352-}
353- 
354-template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>
355-__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyInWeightBias(
356- LocalTensor<float> inputFloat, GlobalTensor<WEIGHT_DTYPE> inputGm, int64_t len)
357-{
358- DataCopyExtParams copyInParams{1, static_cast<uint32_t>(len * sizeof(WEIGHT_DTYPE)), 0, 0, 0};
359- DataCopyPadExtParams<WEIGHT_DTYPE> padParams{false, 0, 0, 0};
360- if constexpr (std::is_same_v<WEIGHT_DTYPE, float>) {
361- DataCopyPad(inputFloat, inputGm, copyInParams, padParams);
362- } else {
363- DataCopyPad(inputFloat.ReinterpretCast<WEIGHT_DTYPE>()[DATA_COUNT], inputGm, copyInParams, padParams);
364- }
365- SetFlag<HardEvent::MTE2_V>(eventIdMte2ToV);
366- WaitFlag<HardEvent::MTE2_V>(eventIdMte2ToV);
367- if constexpr (!std::is_same_v<WEIGHT_DTYPE, float>) {
368- Cast(inputFloat, inputFloat.ReinterpretCast<WEIGHT_DTYPE>()[DATA_COUNT], RoundMode::CAST_NONE, len);
369- PipeBarrier<PIPE_V>();
370- }
371-}
372- 
373template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>333template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>
374__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyInOtherData(int64_t offset, int64_t len)334__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyInOtherData(int64_t offset, int64_t len)
375{335{
376 if (hasWeight) {336 if (hasWeight) {
377- CopyInWeightBias(weightFloat, weightGm[offset], len);337+ CopyInAndCast(weightFloat, weightGm[offset], len, DATA_COUNT);
378 }338 }
379 if (hasBias) {339 if (hasBias) {
380- CopyInWeightBias(biasFloat, biasGm[offset], len);340+ CopyInAndCast(biasFloat, biasGm[offset], len, DATA_COUNT);
381 }341 }
382 if constexpr (OP_CODE == QUANT_OP_CODE) {342 if constexpr (OP_CODE == QUANT_OP_CODE) {
383 if (hasSmooth) {343 if (hasSmooth) {
384- CopyInData(smoothFloat, smoothGm[offset], len);344+ CopyInAndCast(smoothFloat, smoothGm[offset], len, DATA_COUNT);
385 }345 }
386 }346 }
387}347}
@@ -403,8 +363,7 @@ __aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyInSca
403 DataCopyPad(scaleFloat.ReinterpretCast<X_DTYPE>()[DATA_COUNT], scaleGm[offset], copyInParams, padParams);363 DataCopyPad(scaleFloat.ReinterpretCast<X_DTYPE>()[DATA_COUNT], scaleGm[offset], copyInParams, padParams);
404 DataCopyPad(shiftFloat.ReinterpretCast<X_DTYPE>()[DATA_COUNT], shiftGm[offset], copyInParams, padParams);364 DataCopyPad(shiftFloat.ReinterpretCast<X_DTYPE>()[DATA_COUNT], shiftGm[offset], copyInParams, padParams);
405 }365 }
406- SetFlag<HardEvent::MTE2_V>(eventIdMte2ToV);366+ PIPE_MTE2_V();
407- WaitFlag<HardEvent::MTE2_V>(eventIdMte2ToV);
408 int64_t scaleCount = blockCount > 1 ? blockCount * hiddenDimCeil : len;367 int64_t scaleCount = blockCount > 1 ? blockCount * hiddenDimCeil : len;
409 if constexpr (!std::is_same_v<X_DTYPE, float>) {368 if constexpr (!std::is_same_v<X_DTYPE, float>) {
410 Cast(scaleFloat, scaleFloat.ReinterpretCast<X_DTYPE>()[DATA_COUNT], RoundMode::CAST_NONE, scaleCount);369 Cast(scaleFloat, scaleFloat.ReinterpretCast<X_DTYPE>()[DATA_COUNT], RoundMode::CAST_NONE, scaleCount);
@@ -419,29 +378,11 @@ template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>
419__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyInSlice(378__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyInSlice(
420 int64_t offset, int64_t scaleOffset, int64_t h, int64_t len)379 int64_t offset, int64_t scaleOffset, int64_t h, int64_t len)
421{380{
422- CopyInData(xFloat, xGm[offset + h], len);381+ CopyInAndCast(xFloat, xGm[offset + h], len, DATA_COUNT);
423 CopyInScaleShift(scaleOffset, 1, len);382 CopyInScaleShift(scaleOffset, 1, len);
424 CopyInOtherData(h, len);383 CopyInOtherData(h, len);
425}384}
426 385 
427-template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>
428-__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyInSliceX(int64_t offset, int64_t len)
429-{
430- DataCopyExtParams copyInParams{1, static_cast<uint32_t>(len * sizeof(X_DTYPE)), 0, 0, 0};
431- DataCopyPadExtParams<X_DTYPE> padParams{false, 0, 0, 0};
432- if constexpr (std::is_same_v<X_DTYPE, float>) {
433- DataCopyPad(xFloat, xGm[offset], copyInParams, padParams);
434- } else {
435- DataCopyPad(xFloat.ReinterpretCast<X_DTYPE>()[MAX_X_SIZE], xGm[offset], copyInParams, padParams);
436- }
437- SetFlag<HardEvent::MTE2_V>(eventIdMte2ToV);
438- WaitFlag<HardEvent::MTE2_V>(eventIdMte2ToV);
439- if constexpr (!std::is_same_v<X_DTYPE, float>) {
440- Cast(xFloat, xFloat.ReinterpretCast<X_DTYPE>()[MAX_X_SIZE], RoundMode::CAST_NONE, len);
441- PipeBarrier<PIPE_V>();
442- }
443-}
444- 
445template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>386template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>
446__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyInX(int64_t offset, uint16_t blockCount, int64_t len)387__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyInX(int64_t offset, uint16_t blockCount, int64_t len)
447{388{
@@ -457,8 +398,7 @@ __aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyInX(i
457 } else {398 } else {
458 DataCopyPad(xFloat.ReinterpretCast<X_DTYPE>()[DATA_COUNT], xGm[offset], copyInParams, padParams);399 DataCopyPad(xFloat.ReinterpretCast<X_DTYPE>()[DATA_COUNT], xGm[offset], copyInParams, padParams);
459 }400 }
460- SetFlag<HardEvent::MTE2_V>(eventIdMte2ToV);401+ PIPE_MTE2_V();
461- WaitFlag<HardEvent::MTE2_V>(eventIdMte2ToV);
462 if constexpr (!std::is_same_v<X_DTYPE, float>) {402 if constexpr (!std::is_same_v<X_DTYPE, float>) {
463 Cast(xFloat, xFloat.ReinterpretCast<X_DTYPE>()[DATA_COUNT], RoundMode::CAST_NONE, blockCount * hiddenDimCeil);403 Cast(xFloat, xFloat.ReinterpretCast<X_DTYPE>()[DATA_COUNT], RoundMode::CAST_NONE, blockCount * hiddenDimCeil);
464 PipeBarrier<PIPE_V>();404 PipeBarrier<PIPE_V>();
@@ -473,19 +413,7 @@ __aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::BaseCopyO
473 int64_t dataCount = blockCount > 1 ? blockCount * hiddenDimCeil : len;413 int64_t dataCount = blockCount > 1 ? blockCount * hiddenDimCeil : len;
474 Cast(yFloat.ReinterpretCast<X_DTYPE>(), yFloat, RoundMode::CAST_RINT, dataCount);414 Cast(yFloat.ReinterpretCast<X_DTYPE>(), yFloat, RoundMode::CAST_RINT, dataCount);
475 }415 }
476- SetFlag<HardEvent::V_MTE3>(eventIdVToMte3);416+ CopyOut(outGm[offset], yFloat.ReinterpretCast<X_DTYPE>(), blockCount, len);
477- WaitFlag<HardEvent::V_MTE3>(eventIdVToMte3);
478- DataCopyExtParams copyOutParams{blockCount, static_cast<uint32_t>(len * sizeof(X_DTYPE)), 0, 0, 0};
479- DataCopyPad(outGm[offset], yFloat.ReinterpretCast<X_DTYPE>(), copyOutParams);
480-}
481- 
482-template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>
483-__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::QuantCopyOut(int64_t offset, uint16_t blockCount, int64_t len)
484-{
485- SetFlag<HardEvent::V_MTE3>(eventIdVToMte3);
486- WaitFlag<HardEvent::V_MTE3>(eventIdVToMte3);
487- DataCopyExtParams copyOutParams{blockCount, static_cast<uint32_t>(len * sizeof(int8_t)), 0, 0, 0};
488- DataCopyPad(quantOutGm[offset], yInt, copyOutParams);
489}417}
490 418 
491template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>419template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>
@@ -503,6 +431,7 @@ __aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyMeanR
503 if constexpr (!std::is_same_v<X_DTYPE, float>) {431 if constexpr (!std::is_same_v<X_DTYPE, float>) {
504 Cast(meanOutFloat.ReinterpretCast<X_DTYPE>(), meanOutFloat, RoundMode::CAST_RINT, len);432 Cast(meanOutFloat.ReinterpretCast<X_DTYPE>(), meanOutFloat, RoundMode::CAST_RINT, len);
505 Cast(rstdOutFloat.ReinterpretCast<X_DTYPE>(), rstdOutFloat, RoundMode::CAST_RINT, len);433 Cast(rstdOutFloat.ReinterpretCast<X_DTYPE>(), rstdOutFloat, RoundMode::CAST_RINT, len);
434+ event_t eventIdVToMte3 = static_cast<event_t>(pipe.FetchEventID(HardEvent::V_MTE3));
CANN-robot
CANN-robotCANN-robot1月22日

代码可维护性: 在 CopyMeanRstdOut 函数中，动态获取 eventIdVToMte3 事件 ID：event_t eventIdVToMte3 = static_cast<event_t>(pipe.FetchEventID(HardEvent::V_MTE3));。如果 pipe.FetchEventID 返回的值与之前使用的静态事件 ID 不一致，可能导致同步错误。

问题类型: 代码可维护性 文件路径: norm/ada_layer_norm/op_kernel/ada_layer_norm_base_v1.h 行号: 434 问题代码:

event_t eventIdVToMte3 = static_cast<event_t>(pipe.FetchEventID(HardEvent::V_MTE3));

修改建议:

确保动态获取事件 ID 的方式与代码中其他事件同步保持一致。如果其他地方使用了静态事件 ID，建议统一为一种方式，以避免混淆和潜在的错误。

此评论由代码审查工具自动生成

likedislike
506 SetFlag<HardEvent::V_MTE3>(eventIdVToMte3);435 SetFlag<HardEvent::V_MTE3>(eventIdVToMte3);
507 WaitFlag<HardEvent::V_MTE3>(eventIdVToMte3);436 WaitFlag<HardEvent::V_MTE3>(eventIdVToMte3);
508 }437 }
@@ -512,23 +441,3 @@ __aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyMeanR
512 SetFlag<HardEvent::MTE3_S>(eventIdMte3ToS);441 SetFlag<HardEvent::MTE3_S>(eventIdMte3ToS);
513 WaitFlag<HardEvent::MTE3_S>(eventIdMte3ToS);442 WaitFlag<HardEvent::MTE3_S>(eventIdMte3ToS);
514}443}
515- 
516-template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>
517-__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyNormOut(int64_t h, int64_t dataCount)
518-{
519- SetFlag<HardEvent::V_MTE3>(eventIdVToMte3);
520- WaitFlag<HardEvent::V_MTE3>(eventIdVToMte3);
521- DataCopyExtParams copyOutParams{1, static_cast<uint32_t>(dataCount * sizeof(float)), 0, 0, 0};
522- DataCopyPad(normGm[normOffset + h], xFloat, copyOutParams);
523-}
524- 
525-template <typename X_DTYPE, typename WEIGHT_DTYPE, uint8_t OP_CODE>
526-__aicore__ inline void AdaLayerNormND<X_DTYPE, WEIGHT_DTYPE, OP_CODE>::CopyInNorm(int64_t h, int64_t dataCount)
527-{
528- DataCopyExtParams copyInParams{1, static_cast<uint32_t>(dataCount * sizeof(float)), 0, 0, 0};
529- DataCopyPadExtParams<float> padParams{false, 0, 0, 0};
530- DataCopyPad(xFloat, normGm[normOffset + h], copyInParams, padParams);
531- SetFlag<HardEvent::MTE2_V>(eventIdMte2ToV);
532- WaitFlag<HardEvent::MTE2_V>(eventIdMte2ToV);
533-}
534- 
Anorm/ada_layer_norm/op_kernel/ada_layer_norm_util.h+133-0
@@ -0,0 +1,133 @@
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 ada_layer_norm_util.h
13+ * \brief
14+ */
15+#ifndef ADA_LAYER_NORM_UTIL_H
16+#define ADA_LAYER_NORM_UTIL_H
17+ 
18+#include "kernel_operator.h"
19+ 
20+namespace AdaLayerNormNS {
21+using namespace AscendC;
22+ 
23+constexpr uint8_t BASE_OP_CODE = 1;
24+constexpr uint8_t BASE_V2_OP_CODE = 12;
25+constexpr uint8_t QUANT_OP_CODE = 2;
26+constexpr int32_t MAX_X_SIZE = 16384;
27+constexpr int32_t TENSOR_NUM = 7;
28+constexpr int32_t DATA_COUNT = 6144;
29+constexpr int32_t BATCH_COUNT = 1024;
30+constexpr int32_t INT8_BLOCK_NUM = 32;
31+constexpr int32_t HALF_BLOCK_NUM = 16;
32+constexpr int32_t FLOAT_BLOCK_NUM = 8;
33+constexpr float MAX_INT8 = 127.0f;
34+constexpr float ONE_FLOAT = 1.0f;
35+constexpr float FACTOR_INT8 = 1.0f / 127.0f;
36+constexpr int64_t NUM_PER_REP_FP32 = 64;
37+constexpr int64_t MOD_64_MASK = 0x3f;
38+constexpr int64_t LOG2_64 = 6;
39+ 
40+struct RowRange {
41+ int64_t rowStart;
42+ int64_t rowEnd;
43+ int64_t actualRowNum;
44+ int64_t batchStart;
45+ int64_t batchEnd;
46+ int64_t dataCount;
47+};
48+ 
49+struct GmAddr {
50+ const GM_ADDR x = nullptr;
51+ const GM_ADDR scale = nullptr;
52+ const GM_ADDR shift = nullptr;
53+ const GM_ADDR weight = nullptr;
54+ const GM_ADDR bias = nullptr;
55+ const GM_ADDR smooth_scales = nullptr;
56+ const GM_ADDR out = nullptr;
57+ const GM_ADDR mean = nullptr;
58+ const GM_ADDR rstd = nullptr;
59+ const GM_ADDR quant_scale = nullptr;
60+};
61+ 
62+template <typename T1, typename T2>
63+__aicore__ inline T1 CeilA2B(T1 a, T2 b)
64+{
65+ return (b != 0) ? (a + b - 1) / b : a;
66+}
67+ 
68+template <typename T1, typename T2>
69+__aicore__ inline T1 Min(T1 a, T2 b)
70+{
71+ return (a < b) ? a : b;
72+}
73+ 
74+__aicore__ inline void PIPE_MTE2_V()
75+{
76+ event_t eventIDMTE2ToV = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE2_V));
77+ SetFlag<HardEvent::MTE2_V>(eventIDMTE2ToV);
78+ WaitFlag<HardEvent::MTE2_V>(eventIDMTE2ToV);
79+}
80+ 
81+__aicore__ inline void PIPE_V_MTE3()
82+{
83+ event_t eventIdVToMte3 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE3));
84+ SetFlag<HardEvent::V_MTE3>(eventIdVToMte3);
85+ WaitFlag<HardEvent::V_MTE3>(eventIdVToMte3);
86+}
87+ 
88+// count <= 64 * 256(16K)
89+__aicore__ inline void ReduceMaxCustom(const LocalTensor<float>& dstLocal, const LocalTensor<float>& srcLocal, int64_t count)
90+{
91+ int64_t repeatTimes = count >> LOG2_64;
92+ int64_t tailCount = count & MOD_64_MASK;
93+ 
94+ BinaryRepeatParams repeatParams = {1, 1, 1, 0, DEFAULT_REPEAT_STRIDE, 0};
95+ if (likely(repeatTimes > 1)) {
96+ Max(srcLocal, srcLocal[NUM_PER_REP_FP32], srcLocal, NUM_PER_REP_FP32, repeatTimes - 1, repeatParams);
97+ PipeBarrier<PIPE_V>();
98+ }
99+ if (unlikely(tailCount > 0)) {
100+ Max(srcLocal, srcLocal[repeatTimes << LOG2_64], srcLocal, tailCount, 1, repeatParams);
101+ PipeBarrier<PIPE_V>();
102+ }
103+ WholeReduceMax(dstLocal, srcLocal, repeatTimes > 0 ? NUM_PER_REP_FP32 : count, 1, 0, 1, 0, ReduceOrder::ORDER_ONLY_VALUE);
104+}
105+ 
106+template <typename T>
107+__aicore__ inline void CopyOut(const GlobalTensor<T>& dst, const LocalTensor<T>& src, uint16_t blockCount, int64_t len)
108+{
109+ PIPE_V_MTE3();
110+ DataCopyExtParams copyOutParams{blockCount, static_cast<uint32_t>(len * sizeof(T)), 0, 0, 0};
111+ DataCopyPad(dst, src, copyOutParams);
112+}
113+ 
114+template <typename T>
115+__aicore__ inline void CopyInAndCast(const LocalTensor<float>& dst, const GlobalTensor<T>& src, int64_t len, int64_t midOffset)
116+{
117+ DataCopyExtParams copyInParams{1, static_cast<uint32_t>(len * sizeof(T)), 0, 0, 0};
118+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
119+ if constexpr (std::is_same_v<T, float>) {
120+ DataCopyPad(dst, src, copyInParams, padParams);
121+ } else {
122+ DataCopyPad(dst.ReinterpretCast<T>()[midOffset], src, copyInParams, padParams);
123+ }
124+ PIPE_MTE2_V();
125+ if constexpr (!std::is_same_v<T, float>) {
126+ Cast(dst, dst.ReinterpretCast<T>()[midOffset], RoundMode::CAST_NONE, len);
127+ PipeBarrier<PIPE_V>();
128+ }
129+}
130+ 
131+} // namespace AdaLayerNormNS
132+ 
133+#endif // ADA_LAYER_NORM_UTIL_H
Anorm/ada_layer_norm/op_kernel/arch35/ada_layer_norm_common.h+96-0
@@ -0,0 +1,96 @@
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 ada_layer_norm_common.h
13+ * \brief
14+ */
15+#ifndef ADA_LAYER_NORM_COMMON_H
16+#define ADA_LAYER_NORM_COMMON_H
17+ 
18+#include "../ada_layer_norm_util.h"
19+ 
20+namespace AdaLayerNormNS {
21+using namespace AscendC;
22+using namespace AscendC::MicroAPI;
23+ 
24+constexpr int32_t DOUBLE_BUFFER = 2;
25+constexpr uint16_t V_LENGTH = VECTOR_REG_WIDTH / sizeof(float);
26+constexpr uint16_t TWO_V_LENGTH = V_LENGTH << 1;
27+constexpr float FACTOR_FP8_E5M2 = 1.0f / 57344.0f;
28+constexpr float FACTOR_FP8_E4M3FN = 1.0f / 448.0f;
29+constexpr float FACTOR_HIFLOAT8 = 1.0f / 32768.0f;
30+ 
31+constexpr CastTrait castTraitB16ToB32 = {
32+ RegLayout::ZERO, SatMode::UNKNOWN, MaskMergeMode::ZEROING, RoundMode::UNKNOWN};
33+constexpr CastTrait castTraitB32ToB16 = {
34+ RegLayout::ZERO, SatMode::NO_SAT, MaskMergeMode::ZEROING, RoundMode::CAST_RINT};
35+constexpr CastTrait castTraitF32ToI16 = {
36+ RegLayout::ZERO, SatMode::NO_SAT, MaskMergeMode::ZEROING, RoundMode::CAST_RINT};
37+constexpr CastTrait castTraitI16ToF16 = {
38+ RegLayout::ZERO, SatMode::UNKNOWN, MaskMergeMode::ZEROING, RoundMode::CAST_ROUND};
39+constexpr CastTrait castTraitF16ToI8 = {
40+ RegLayout::ZERO, SatMode::NO_SAT, MaskMergeMode::ZEROING, RoundMode::CAST_TRUNC};
41+constexpr CastTrait castTraitF32Tofp8 = {
42+ RegLayout::ZERO, SatMode::NO_SAT, MaskMergeMode::ZEROING, RoundMode::CAST_RINT};
43+constexpr CastTrait castTraitF32Toh8 = {
44+ RegLayout::ZERO, SatMode::NO_SAT, MaskMergeMode::ZEROING, RoundMode::CAST_ROUND};
45+ 
46+constexpr LayerNormConfig hasGammaBetaConfig = {false, false, false};
47+constexpr LayerNormConfig hasGammaNoBetaConfig = {true, false, false};
48+constexpr LayerNormConfig noGammaHasBetaConfig = {false, true, false};
49+constexpr LayerNormConfig noGammaNoBetaConfig = {true, true, false};
CANN-robot
CANN-robotCANN-robot1月22日

代码结构与可维护性: noGammaNoBetaConfig常量的初始化值与命名存在矛盾。常量名'noGammaNoBetaConfig'暗示没有Gamma和Beta，但其初始化值却为{true, true, false}，前两个布尔值（通常对应hasGamma和hasBeta）被设置为true，这与命名相悖。从上下文看，第46-49行定义了四个配置，分别对应Gamma和Beta的有无组合。根据命名惯例，'hasGammaBetaConfig'表示两者都有（false, false, false？这里需要查看LayerNormConfig结构定义），但当前初始化值逻辑不清晰，容易导致使用错误。

问题类型: 代码结构与可维护性 文件路径: norm/ada_layer_norm/op_kernel/arch35/ada_layer_norm_common.h 行号: 49 问题代码:

constexpr LayerNormConfig noGammaNoBetaConfig = {true, true, false};

修改建议:

1. 检查LayerNormConfig结构体的定义，明确各布尔成员的含义（例如，是否是hasGamma, hasBeta, 或其他）。
2. 根据结构体成员的实际含义，修正noGammaNoBetaConfig的初始化值，使其与命名一致。例如，如果前两个成员分别表示hasGamma和hasBeta，则应改为{false, false, false}。
3. 同样，检查并修正其他三个配置常量（hasGammaBetaConfig, hasGammaNoBetaConfig, noGammaHasBetaConfig）的初始化值，确保其与命名匹配。
4. 在常量定义处添加注释，说明每个布尔值的具体含义。

此评论由代码审查工具自动生成

likedislike
50+ 
51+constexpr NormalizeConfig hasGammaBetaNormalizeConfig = {ReducePattern::AR, -1, false, false, false};
52+constexpr NormalizeConfig hasGammaNoBetaNormalizeConfig = {ReducePattern::AR, -1, true, false, false};
53+constexpr NormalizeConfig noGammaHasBetaNormalizeConfig = {ReducePattern::AR, -1, false, true, false};
54+constexpr NormalizeConfig noGammaNoBetaNormalizeConfig = {ReducePattern::AR, -1, true, true, false};
55+ 
56+template <typename T>
57+__aicore__ inline void LoadTensor(RegTensor<float>& dst, __local_mem__ T* srcAddr, MaskReg& pregLoop)
58+{
59+ if constexpr (std::is_same_v<T, float>) {
60+ DataCopy(dst, srcAddr);
61+ } else {
62+ RegTensor<T> tmpFp16;
63+ DataCopy<T, LoadDist::DIST_UNPACK_B16>(tmpFp16, srcAddr);
64+ Cast<float, T, castTraitB16ToB32>(dst, tmpFp16, pregLoop);
65+ }
66+}
67+ 
68+template <typename T>
69+__aicore__ inline void CopyToTensor(__local_mem__ T* dstAddr, RegTensor<float>& src, MaskReg& pregLoop)
70+{
71+ if constexpr (std::is_same_v<T, float>) {
72+ DataCopy(dstAddr, src, pregLoop);
73+ } else {
74+ RegTensor<T> tmpFp16;
75+ Cast<T, float, castTraitB32ToB16>(tmpFp16, src, pregLoop);
76+ DataCopy<T, StoreDist::DIST_PACK_B32>(dstAddr, tmpFp16, pregLoop);
77+ }
78+}
79+ 
80+template <typename T>
81+__aicore__ inline float GetQuantFactor()
82+{
83+ if constexpr (std::is_same_v<T, hifloat8_t>) {
84+ return FACTOR_HIFLOAT8;
85+ } else if constexpr (std::is_same_v<T, fp8_e4m3fn_t>) {
86+ return FACTOR_FP8_E4M3FN;
87+ } else if constexpr (std::is_same_v<T, fp8_e5m2_t>) {
88+ return FACTOR_FP8_E5M2;
89+ } else {
90+ return FACTOR_INT8;
91+ }
92+}
93+ 
94+} // namespace AdaLayerNormNS
95+ 
96+#endif // ADA_LAYER_NORM_COMMON_H
Anorm/ada_layer_norm/op_kernel/arch35/ada_layer_norm_full_load.h+495-0
@@ -0,0 +1,495 @@
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 ada_layer_norm_full_load.h
13+ * \brief
14+ */
15+#ifndef ADA_LAYER_NORM_FULL_LOAD_H
16+#define ADA_LAYER_NORM_FULL_LOAD_H
17+ 
18+#include "ada_layer_norm_common.h"
19+ 
20+namespace AdaLayerNormNS {
21+using namespace AscendC;
22+ 
23+template <typename T, typename U, typename Y, uint8_t OP_CODE>
24+class AdaLayerNormFullLoad {
25+public:
26+ __aicore__ inline AdaLayerNormFullLoad(){};
27+ __aicore__ inline void InitV2(const GmAddr* gmAddr, const AdaLayerNormTilingData *tilingData);
28+ __aicore__ inline void InitQuant(const GmAddr* gmAddr, GM_ADDR workspace, const AdaLayerNormTilingData *tilingData);
29+ __aicore__ inline void Process();
30+ 
31+private:
32+ __aicore__ inline void ParseTilingData(const AdaLayerNormTilingData *tilingData);
33+ __aicore__ inline void FastProcess();
34+ __aicore__ inline void ProcessLayerNorm(RowRange range, int64_t batchCount);
35+ __aicore__ inline void Adaption(RowRange range, int64_t batchCount);
36+ 
37+ __aicore__ inline void CopyInOtherData();
38+ __aicore__ inline void CopyInScaleShift(int64_t offset, uint16_t blockCount, int64_t len);
39+ __aicore__ inline void CopyInX(int64_t offset, uint16_t blockCount, int64_t len);
40+ __aicore__ inline void BaseCopyOut(int64_t offset, uint16_t blockCount, int64_t len);
41+ __aicore__ inline void QuantCopyOut(int64_t offset, uint16_t blockCount, int64_t len);
42+ __aicore__ inline void CopyMeanRstdOut(int64_t offset, int64_t len);
43+ __aicore__ inline void CopyScaleOut(int64_t offset, int64_t len);
44+ 
45+private:
46+ using OUT_DTYPE = std::conditional_t<OP_CODE == QUANT_OP_CODE, float, T>;
47+ TPipe pipe;
48+ RowRange range;
49+ TQue<QuePosition::VECIN, DOUBLE_BUFFER> xQueue;
50+ TQue<QuePosition::VECIN, 1> scaleQueue;
51+ TQue<QuePosition::VECIN, 1> shiftQueue;
52+ TQue<QuePosition::VECIN, 1> smoothQueue;
53+ TBuf<TPosition::VECCALC> weightBuf;
54+ TBuf<TPosition::VECCALC> biasBuf;
55+ TBuf<TPosition::VECCALC> normBuf;
56+ TBuf<TPosition::VECCALC> tmpBuf;
57+ TQue<QuePosition::VECOUT, DOUBLE_BUFFER> outQueue;
58+ TQue<QuePosition::VECOUT, DOUBLE_BUFFER> quantOutQueue;
59+ TQue<QuePosition::VECOUT, 1> meanQueue;
60+ TQue<QuePosition::VECOUT, 1> rstdQueue;
61+ TQue<QuePosition::VECOUT, 1> quantScaleQueue;
62+ 
63+ GlobalTensor<T> xGm;
64+ GlobalTensor<T> scaleGm;
65+ GlobalTensor<T> shiftGm;
66+ GlobalTensor<T> smoothGm;
67+ GlobalTensor<U> weightGm;
68+ GlobalTensor<U> biasGm;
69+ GlobalTensor<T> outGm;
70+ GlobalTensor<T> meanGm;
71+ GlobalTensor<T> rstdGm;
72+ GlobalTensor<Y> quantOutGm;
73+ GlobalTensor<float> quantScaleGm;
74+ 
75+ LocalTensor<float> weightLocal;
76+ LocalTensor<float> biasLocal;
77+ LocalTensor<T> smoothLocal;
78+ LocalTensor<T> scaleLocal;
79+ LocalTensor<T> shiftLocal;
80+ LocalTensor<float> normLocal;
81+ LocalTensor<uint8_t> tmpLocal;
82+ LocalTensor<float> meanLocal;
83+ LocalTensor<float> rstdLocal;
84+ LocalTensor<float> quantScaleLocal;
85+ 
86+ int64_t rowNum = 0;
87+ int64_t seqLen = 0;
88+ int64_t hiddenDim = 0;
89+ int64_t hiddenDimCeil = 0;
90+ int64_t outIntCeil = 0;
91+ int64_t rowStart = 0;
92+ int64_t rowEnd = 0;
93+ int64_t tmpBufferSize = 0;
94+ float epsilon = 0.0f;
95+ int32_t hasWeight = 0;
96+ int32_t hasBias = 0;
97+ int32_t hasSmooth = 0;
98+ float quantFactor = 0.0f;
99+ LayerNormSeparateTiling layerNormTiling;
100+};
101+ 
102+template <typename T, typename U, typename Y, uint8_t OP_CODE>
103+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::InitV2(const GmAddr* gmAddr, const AdaLayerNormTilingData *tilingData)
104+{
105+ ParseTilingData(tilingData);
106+ pipe.InitBuffer(xQueue, DOUBLE_BUFFER, DATA_COUNT * sizeof(float));
107+ pipe.InitBuffer(scaleQueue, 1, DATA_COUNT * sizeof(float));
108+ pipe.InitBuffer(shiftQueue, 1, DATA_COUNT * sizeof(float));
109+ pipe.InitBuffer(weightBuf, DATA_COUNT * sizeof(float));
110+ pipe.InitBuffer(biasBuf, DATA_COUNT * sizeof(float));
111+ pipe.InitBuffer(normBuf, DATA_COUNT * sizeof(float));
112+ if (tmpBufferSize > 0) {
113+ pipe.InitBuffer(tmpBuf, tmpBufferSize);
114+ }
115+ pipe.InitBuffer(outQueue, DOUBLE_BUFFER, DATA_COUNT * sizeof(float));
116+ pipe.InitBuffer(meanQueue, DOUBLE_BUFFER, BATCH_COUNT * sizeof(float));
117+ pipe.InitBuffer(rstdQueue, DOUBLE_BUFFER, BATCH_COUNT * sizeof(float));
118+ 
119+ xGm.SetGlobalBuffer((__gm__ T*)gmAddr->x);
120+ weightGm.SetGlobalBuffer((__gm__ U*)gmAddr->weight);
121+ biasGm.SetGlobalBuffer((__gm__ U*)gmAddr->bias);
122+ scaleGm.SetGlobalBuffer((__gm__ T*)gmAddr->scale);
123+ shiftGm.SetGlobalBuffer((__gm__ T*)gmAddr->shift);
124+ outGm.SetGlobalBuffer((__gm__ T*)gmAddr->out);
125+ meanGm.SetGlobalBuffer((__gm__ T*)gmAddr->mean);
126+ rstdGm.SetGlobalBuffer((__gm__ T*)gmAddr->rstd);
127+}
128+ 
129+template <typename T, typename U, typename Y, uint8_t OP_CODE>
130+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::InitQuant(const GmAddr* gmAddr, GM_ADDR workspace,
131+ const AdaLayerNormTilingData *tilingData)
132+{
133+ ParseTilingData(tilingData);
134+ pipe.InitBuffer(xQueue, DOUBLE_BUFFER, DATA_COUNT * sizeof(float));
135+ pipe.InitBuffer(scaleQueue, 1, DATA_COUNT * sizeof(float));
136+ pipe.InitBuffer(shiftQueue, 1, DATA_COUNT * sizeof(float));
137+ pipe.InitBuffer(smoothQueue, 1, DATA_COUNT * sizeof(float));
138+ pipe.InitBuffer(weightBuf, DATA_COUNT * sizeof(float));
139+ pipe.InitBuffer(biasBuf, DATA_COUNT * sizeof(float));
140+ pipe.InitBuffer(normBuf, DATA_COUNT * sizeof(float));
141+ if (tmpBufferSize > 0) {
142+ pipe.InitBuffer(tmpBuf, tmpBufferSize);
143+ }
144+ pipe.InitBuffer(outQueue, 1, DATA_COUNT * sizeof(float));
145+ pipe.InitBuffer(quantOutQueue, DOUBLE_BUFFER, DATA_COUNT * sizeof(int8_t));
146+ pipe.InitBuffer(meanQueue, 1, BATCH_COUNT * sizeof(float));
147+ pipe.InitBuffer(rstdQueue, 1, BATCH_COUNT * sizeof(float));
148+ pipe.InitBuffer(quantScaleQueue, 1, BATCH_COUNT * sizeof(float));
149+ 
150+ xGm.SetGlobalBuffer((__gm__ T*)gmAddr->x);
151+ scaleGm.SetGlobalBuffer((__gm__ T*)gmAddr->scale);
152+ shiftGm.SetGlobalBuffer((__gm__ T*)gmAddr->shift);
153+ weightGm.SetGlobalBuffer((__gm__ U*)gmAddr->weight);
154+ biasGm.SetGlobalBuffer((__gm__ U*)gmAddr->bias);
155+ smoothGm.SetGlobalBuffer((__gm__ T*)gmAddr->smooth_scales);
156+ quantOutGm.SetGlobalBuffer((__gm__ Y*)gmAddr->out);
157+ quantScaleGm.SetGlobalBuffer((__gm__ float*)gmAddr->quant_scale);
158+}
159+ 
160+template <typename T, typename U, typename Y, uint8_t OP_CODE>
161+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::Process()
162+{
163+ if (rowStart >= rowEnd) {
164+ return;
165+ }
166+ 
167+ weightLocal = weightBuf.Get<float>();
168+ biasLocal = biasBuf.Get<float>();
169+ normLocal = normBuf.Get<float>();
170+ tmpLocal = tmpBuf.Get<uint8_t>();
171+ meanLocal = meanQueue.AllocTensor<float>();
172+ rstdLocal = rstdQueue.AllocTensor<float>();
173+ CopyInOtherData();
174+ if constexpr (OP_CODE == QUANT_OP_CODE) {
175+ if (hasSmooth) {
176+ smoothLocal = smoothQueue.DeQue<T>();
177+ }
178+ quantScaleLocal = quantScaleQueue.AllocTensor<float>();
179+ }
180+ FastProcess();
181+ scaleQueue.FreeTensor(scaleLocal);
182+ shiftQueue.FreeTensor(shiftLocal);
183+ if constexpr (OP_CODE == QUANT_OP_CODE) {
184+ if (hasSmooth) {
185+ smoothQueue.FreeTensor(smoothLocal);
186+ }
187+ }
188+}
189+ 
190+template <typename T, typename U, typename Y, uint8_t OP_CODE>
191+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::ParseTilingData(const AdaLayerNormTilingData *tilingData)
192+{
193+ int32_t blockNum = HALF_BLOCK_NUM;
194+ if constexpr (std::is_same_v<T, float>) {
195+ blockNum = FLOAT_BLOCK_NUM;
196+ }
197+ seqLen = tilingData->seqLen;
198+ hiddenDim = tilingData->hiddenDim;
199+ hiddenDimCeil = CeilA2B(hiddenDim, blockNum) * blockNum;
200+ outIntCeil = CeilA2B(hiddenDim, INT8_BLOCK_NUM) * INT8_BLOCK_NUM;
201+ epsilon = tilingData->epsilon;
202+ hasWeight = tilingData->hasWeight;
203+ hasBias = tilingData->hasBias;
204+ hasSmooth = tilingData->hasSmooth;
205+ rowNum = tilingData->rowNum;
206+ tmpBufferSize = tilingData->tmpBufferSize;
207+ layerNormTiling = tilingData->layerNormTiling;
208+ quantFactor = GetQuantFactor<Y>();
209+ 
210+ int64_t singleCoreNum = tilingData->singleCoreNum;
211+ int64_t tailNum = tilingData->tailNum;
212+ int64_t blockIdx = GetBlockIdx();
213+ rowStart = singleCoreNum * blockIdx + (blockIdx < tailNum ? blockIdx : tailNum);
214+ rowEnd = rowStart + singleCoreNum + (blockIdx < tailNum ? 1 : 0);
215+}
216+ 
217+template <typename T, typename U, typename Y, uint8_t OP_CODE>
218+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::FastProcess()
219+{
220+ int64_t lastBatchStart = -1;
221+ int64_t lastBatchEnd = -1;
222+ int64_t rowIdx = rowStart;
223+ int64_t batchCount = 0;
224+ while (rowIdx < rowEnd) {
225+ range.rowStart = rowIdx;
226+ range.rowEnd = Min(rowEnd, range.rowStart + rowNum);
227+ range.batchStart = range.rowStart / seqLen;
228+ range.batchEnd = (range.rowEnd - 1) / seqLen + 1;
229+ range.actualRowNum = range.rowEnd - range.rowStart;
230+ range.dataCount = range.actualRowNum > 1 ? range.actualRowNum * hiddenDimCeil : hiddenDim;
231+ 
232+ // 拷入
233+ bool scaleShiftReuse = (lastBatchStart == range.batchStart && lastBatchEnd == range.batchEnd);
234+ if (!scaleShiftReuse) {
235+ if (lastBatchStart != -1) {
236+ scaleQueue.FreeTensor(scaleLocal);
237+ shiftQueue.FreeTensor(shiftLocal);
238+ }
239+ CopyInScaleShift(range.batchStart * hiddenDim, range.batchEnd - range.batchStart, hiddenDim);
240+ scaleLocal = scaleQueue.DeQue<T>();
241+ shiftLocal = shiftQueue.DeQue<T>();
242+ lastBatchStart = range.batchStart;
243+ lastBatchEnd = range.batchEnd;
244+ }
245+ CopyInX(range.rowStart * hiddenDim, range.actualRowNum, hiddenDim);
246+ 
247+ // 计算 & 拷出
248+ ProcessLayerNorm(range, batchCount);
249+ Adaption(range, batchCount);
250+ if constexpr (OP_CODE == QUANT_OP_CODE) {
251+ QuantCopyOut(range.rowStart * hiddenDim, range.actualRowNum, hiddenDim);
252+ } else {
253+ BaseCopyOut(range.rowStart * hiddenDim, range.actualRowNum, hiddenDim);
254+ }
255+ 
256+ batchCount += range.actualRowNum;
257+ rowIdx = range.rowEnd;
258+ if (batchCount + rowNum > BATCH_COUNT || rowIdx == rowEnd) {
259+ if constexpr (OP_CODE == QUANT_OP_CODE) {
260+ quantScaleQueue.EnQue<float>(quantScaleLocal);
261+ CopyScaleOut(rowIdx - batchCount, batchCount);
262+ if (rowIdx != rowEnd) {
263+ quantScaleLocal = quantScaleQueue.AllocTensor<float>();
264+ }
265+ } else {
266+ if constexpr (std::is_same_v<T, float>) {
267+ meanQueue.EnQue<float>(meanLocal);
268+ rstdQueue.EnQue<float>(rstdLocal);
269+ } else {
270+ LocalTensor<T> meanOut = meanLocal.ReinterpretCast<T>();
271+ LocalTensor<T> rstdOut = rstdLocal.ReinterpretCast<T>();
272+ Cast(meanOut, meanLocal, RoundMode::CAST_RINT, batchCount);
273+ Cast(rstdOut, rstdLocal, RoundMode::CAST_RINT, batchCount);
274+ meanQueue.EnQue<T>(meanOut);
275+ rstdQueue.EnQue<T>(rstdOut);
276+ }
277+ CopyMeanRstdOut(rowIdx - batchCount, batchCount);
278+ if (rowIdx != rowEnd) {
279+ meanLocal = meanQueue.AllocTensor<float>();
280+ rstdLocal = rstdQueue.AllocTensor<float>();
281+ }
282+ }
283+ batchCount = 0;
284+ }
285+ }
286+}
287+ 
288+template <typename T, typename U, typename Y, uint8_t OP_CODE>
289+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::ProcessLayerNorm(RowRange range, int64_t batchCount)
290+{
291+ LocalTensor<float> xLocal;
292+ if constexpr (std::is_same_v<T, float>) {
293+ xLocal = xQueue.DeQue<float>();
294+ } else {
295+ LocalTensor<T> xFp16 = xQueue.DeQue<T>();
296+ xLocal = xFp16.template ReinterpretCast<float>();
297+ Cast(xLocal, xFp16[DATA_COUNT], RoundMode::CAST_NONE, range.dataCount);
298+ PipeBarrier<PIPE_V>();
299+ }
300+ 
301+ LayerNormPara para;
302+ para.aLength = range.actualRowNum;
303+ para.rLength = hiddenDim;
304+ para.rLengthWithPadding = hiddenDimCeil;
305+ if (hasWeight && hasBias) {
306+ LayerNorm<float, float, true, hasGammaBetaConfig>(
307+ normLocal, meanLocal[batchCount], rstdLocal[batchCount], xLocal, weightLocal, biasLocal,
308+ epsilon, tmpLocal, para, layerNormTiling);
309+ } else if (!hasWeight && hasBias) {
310+ LayerNorm<float, float, true, noGammaHasBetaConfig>(
311+ normLocal, meanLocal[batchCount], rstdLocal[batchCount], xLocal, weightLocal, biasLocal,
312+ epsilon, tmpLocal, para, layerNormTiling);
313+ } else if (hasWeight && !hasBias) {
314+ LayerNorm<float, float, true, hasGammaNoBetaConfig>(
315+ normLocal, meanLocal[batchCount], rstdLocal[batchCount], xLocal, weightLocal, biasLocal,
316+ epsilon, tmpLocal, para, layerNormTiling);
317+ } else {
318+ LayerNorm<float, float, true, noGammaNoBetaConfig>(
319+ normLocal, meanLocal[batchCount], rstdLocal[batchCount], xLocal, weightLocal, biasLocal,
320+ epsilon, tmpLocal, para, layerNormTiling);
321+ }
322+ PipeBarrier<PIPE_V>();
323+ xQueue.FreeTensor(xLocal);
324+}
325+ 
326+template <typename T, typename U, typename Y, uint8_t OP_CODE>
327+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::Adaption(RowRange range, int64_t batchCount)
328+{
329+ LocalTensor<OUT_DTYPE> outLocal = outQueue.AllocTensor<OUT_DTYPE>();
330+ __local_mem__ float* normAddr = (__ubuf__ float*)normLocal.GetPhyAddr();
331+ __local_mem__ OUT_DTYPE* outAddr = (__ubuf__ OUT_DTYPE*)outLocal.GetPhyAddr();
332+ __local_mem__ T* scaleAddr = (__ubuf__ T*)scaleLocal.GetPhyAddr();
333+ __local_mem__ T* shiftAddr = (__ubuf__ T*)shiftLocal.GetPhyAddr();
334+ 
335+ LocalTensor<Y> quantOutLocal;
336+ __local_mem__ Y* quantOutAddr;
337+ __local_mem__ T* smoothAddr;
338+ __local_mem__ float* quantScaleAddr;
339+ if constexpr (OP_CODE == QUANT_OP_CODE) {
340+ quantOutLocal = quantOutQueue.AllocTensor<Y>();
341+ quantOutAddr = (__ubuf__ Y*)quantOutLocal.GetPhyAddr();
342+ smoothAddr = (__ubuf__ T*)smoothLocal.GetPhyAddr();
CANN-robot
CANN-robotCANN-robot1月22日

指针与引用安全: 在Adaption函数中，当OP_CODE == QUANT_OP_CODE时，smoothAddr被赋值为(ubuf T*)smoothLocal.GetPhyAddr()。但smoothLocal仅在hasSmooth为真时才从smoothQueue.DeQue()获取（见Process函数第176行）。如果hasSmooth为假，smoothLocal可能未初始化或无效，而smoothAddr仍被使用（第399-403行），导致未定义行为。

问题类型: 指针与引用安全 文件路径: norm/ada_layer_norm/op_kernel/arch35/ada_layer_norm_full_load.h 行号: 342 问题代码:

smoothAddr = (__ubuf__ T*)smoothLocal.GetPhyAddr();

修改建议:

将smoothAddr的赋值和smoothLocal的使用包裹在if (hasSmooth)条件内，或者确保当hasSmooth为假时，smoothAddr不被使用。

此评论由代码审查工具自动生成

likedislike
343+ quantScaleAddr = (__ubuf__ float*)quantScaleLocal[batchCount].GetPhyAddr();
344+ }
345+ 
346+ uint16_t rowLoopTimes = static_cast<uint16_t>(range.actualRowNum);
347+ uint16_t tailLength = static_cast<uint16_t>(hiddenDim % TWO_V_LENGTH);
348+ uint16_t colLoopTimes = static_cast<uint16_t>(hiddenDim / TWO_V_LENGTH) + (tailLength > V_LENGTH ? 1 : 0);
349+ uint32_t rightLength = static_cast<uint32_t>(hiddenDim - colLoopTimes * V_LENGTH);
350+ __VEC_SCOPE__
351+ {
352+ RegTensor<float> x1;
353+ RegTensor<float> x2;
354+ RegTensor<float> scale1;
355+ RegTensor<float> scale2;
356+ RegTensor<float> shift1;
357+ RegTensor<float> shift2;
358+ RegTensor<float> tmpMax;
359+ RegTensor<float> quantScale;
360+ RegTensor<float> quantScaleBroad;
361+ RegTensor<int16_t> y1Int16;
362+ RegTensor<int16_t> y2Int16;
363+ RegTensor<half> y1Fp16;
364+ RegTensor<half> y2Fp16;
365+ RegTensor<Y> y1;
366+ RegTensor<Y> y2;
367+ 
368+ MaskReg pregFull = CreateMask<float, MaskPattern::ALL>();
369+ MaskReg pregMerge = CreateMask<float, MaskPattern::VL1>();
370+ MaskReg pregLoop;
371+ __local_mem__ OUT_DTYPE* outAddr1 = outAddr;
372+ __local_mem__ OUT_DTYPE* outAddr2 = outAddr;
373+ for (uint16_t i = 0; i < rowLoopTimes; i++) {
374+ uint32_t batchIdx = (static_cast<uint32_t>(range.rowStart) + i) / static_cast<uint32_t>(seqLen);
375+ uint32_t scaleOffset = (batchIdx - static_cast<uint32_t>(range.batchStart)) * static_cast<uint32_t>(hiddenDimCeil);
376+ uint32_t sreg1 = rightLength;
377+ if constexpr (OP_CODE == QUANT_OP_CODE) {
378+ Duplicate(tmpMax, 0.0f, pregFull);
379+ }
380+ for (uint16_t j = 0; j < colLoopTimes;j ++) {
381+ pregLoop = UpdateMask<float>(sreg1);
382+ DataCopy(x1, normAddr + j * TWO_V_LENGTH);
383+ DataCopy(x2, normAddr + j * TWO_V_LENGTH + V_LENGTH);
384+ LoadTensor(scale1, scaleAddr + scaleOffset + j * TWO_V_LENGTH, pregFull);
385+ LoadTensor(scale2, scaleAddr + scaleOffset + j * TWO_V_LENGTH + V_LENGTH, pregLoop);
386+ LoadTensor(shift1, shiftAddr + scaleOffset + j * TWO_V_LENGTH, pregFull);
387+ LoadTensor(shift2, shiftAddr + scaleOffset + j * TWO_V_LENGTH + V_LENGTH, pregLoop);
388+ Adds(scale1, scale1, 1.0f, pregFull);
389+ Adds(scale2, scale2, 1.0f, pregLoop);
390+ FusedMulDstAdd(x1, scale1, shift1, pregFull);
391+ FusedMulDstAdd(x2, scale2, shift2, pregLoop);
392+ if constexpr (OP_CODE == BASE_V2_OP_CODE) {
393+ CopyToTensor(outAddr1 + j * TWO_V_LENGTH, x1, pregFull);
394+ CopyToTensor(outAddr1 + j * TWO_V_LENGTH + V_LENGTH, x2, pregLoop);
395+ } else {
396+ if (hasSmooth) {
397+ RegTensor<float> smooth1;
398+ RegTensor<float> smooth2;
399+ LoadTensor(smooth1, smoothAddr + j * TWO_V_LENGTH, pregFull);
400+ LoadTensor(smooth2, smoothAddr + j * TWO_V_LENGTH + V_LENGTH, pregLoop);
401+ Mul(x1, x1, smooth1, pregFull);
402+ Mul(x2, x2, smooth2, pregLoop);
403+ }
404+ CopyToTensor(outAddr1 + j * TWO_V_LENGTH, x1, pregFull);
405+ CopyToTensor(outAddr1 + j * TWO_V_LENGTH + V_LENGTH, x2, pregLoop);
406+ Abs(x1, x1, pregFull);
407+ Abs(x2, x2, pregLoop);
408+ Max(x1, x1, x2, pregFull);
409+ Max(tmpMax, tmpMax, x1, pregFull);
410+ }
411+ }
412+ if (tailLength > 0 && tailLength <= V_LENGTH) {
413+ pregLoop = UpdateMask<float>(sreg1);
414+ DataCopy(x1, normAddr + colLoopTimes * TWO_V_LENGTH);
415+ LoadTensor(scale1, scaleAddr + scaleOffset + colLoopTimes * TWO_V_LENGTH, pregLoop);
416+ LoadTensor(shift1, shiftAddr + scaleOffset + colLoopTimes * TWO_V_LENGTH, pregLoop);
417+ Adds(scale1, scale1, 1.0f, pregLoop);
418+ FusedMulDstAdd(x1, scale1, shift1, pregLoop);
419+ if constexpr (OP_CODE == QUANT_OP_CODE) {
420+ if (hasSmooth) {
421+ RegTensor<float> smooth1;
422+ LoadTensor(smooth1, smoothAddr + colLoopTimes * TWO_V_LENGTH, pregLoop);
423+ Mul(x1, x1, smooth1, pregLoop);
424+ }
425+ CopyToTensor(outAddr1 + colLoopTimes * TWO_V_LENGTH, x1, pregLoop);
426+ Abs(x1, x1, pregLoop);
427+ Max(tmpMax, tmpMax, x1, pregFull);
428+ } else {
429+ CopyToTensor(outAddr1 + colLoopTimes * TWO_V_LENGTH, x1, pregLoop);
430+ }
431+ }
432+ if constexpr (OP_CODE == QUANT_OP_CODE) {
433+ ReduceMax(quantScale, tmpMax, pregFull);
434+ Muls(quantScale, quantScale, quantFactor, pregMerge);
435+ Duplicate(quantScaleBroad, quantScale, pregFull);
436+ LocalMemBar<MemType::VEC_STORE, MemType::VEC_LOAD>();
437+ uint32_t sreg2 = rightLength;
438+ for (uint16_t j = 0; j < colLoopTimes;j ++) {
439+ pregLoop = UpdateMask<float>(sreg2);
440+ DataCopy(x1, outAddr2 + j * TWO_V_LENGTH);
441+ DataCopy(x2, outAddr2 + j * TWO_V_LENGTH + V_LENGTH);
442+ Div(x1, x1, quantScaleBroad, pregFull);
443+ Div(x2, x2, quantScaleBroad, pregLoop);
444+ if constexpr (std::is_same_v<Y, hifloat8_t>) {
445+ Cast<Y, float, castTraitF32Toh8>(y1, x1, pregFull);
446+ Cast<Y, float, castTraitF32Toh8>(y2, x2, pregLoop);
447+ } else if constexpr (std::is_same_v<Y, int8_t>) {
448+ Cast<int16_t, float, castTraitF32ToI16>(y1Int16, x1, pregFull);
449+ Cast<int16_t, float, castTraitF32ToI16>(y2Int16, x2, pregLoop);
450+ Cast<half, int16_t, castTraitI16ToF16>(y1Fp16, y1Int16, pregFull);
451+ Cast<half, int16_t, castTraitI16ToF16>(y2Fp16, y2Int16, pregLoop);
452+ Cast<Y, half, castTraitF16ToI8>(y1, y1Fp16, pregFull);
453+ Cast<Y, half, castTraitF16ToI8>(y2, y2Fp16, pregLoop);
454+ } else {
455+ Cast<Y, float, castTraitF32Tofp8>(y1, x1, pregFull);
456+ Cast<Y, float, castTraitF32Tofp8>(y2, x2, pregLoop);
457+ }
458+ DataCopy<Y, StoreDist::DIST_PACK4_B32>(quantOutAddr + j * TWO_V_LENGTH, y1, pregFull);
459+ DataCopy<Y, StoreDist::DIST_PACK4_B32>(quantOutAddr + j * TWO_V_LENGTH + V_LENGTH, y2, pregLoop);
460+ }
461+ if (tailLength > 0 && tailLength <= V_LENGTH) {
462+ pregLoop = UpdateMask<float>(sreg2);
463+ DataCopy(x1, outAddr2 + colLoopTimes * TWO_V_LENGTH);
464+ Div(x1, x1, quantScaleBroad, pregLoop);
465+ if constexpr (std::is_same_v<Y, hifloat8_t>) {
466+ Cast<Y, float, castTraitF32Toh8>(y1, x1, pregLoop);
467+ } else if constexpr (std::is_same_v<Y, int8_t>) {
468+ Cast<int16_t, float, castTraitF32ToI16>(y1Int16, x1, pregLoop);
469+ Cast<half, int16_t, castTraitI16ToF16>(y1Fp16, y1Int16, pregLoop);
470+ Cast<Y, half, castTraitF16ToI8>(y1, y1Fp16, pregLoop);
471+ } else {
472+ Cast<Y, float, castTraitF32Tofp8>(y1, x1, pregLoop);
473+ }
474+ DataCopy<Y, StoreDist::DIST_PACK4_B32>(quantOutAddr + colLoopTimes * TWO_V_LENGTH, y1, pregLoop);
475+ }
476+ // 拷出量化系数
477+ DataCopy<float, StoreDist::DIST_FIRST_ELEMENT_B32>(quantScaleAddr + i, quantScale, pregMerge);
478+ 
479+ outAddr2 += static_cast<uint32_t>(hiddenDimCeil);
480+ quantOutAddr += static_cast<uint32_t>(outIntCeil);
481+ }
482+ normAddr += static_cast<uint32_t>(hiddenDimCeil);
483+ outAddr1 += static_cast<uint32_t>(hiddenDimCeil);
484+ }
485+ }
486+ if constexpr (OP_CODE == QUANT_OP_CODE) {
487+ quantOutQueue.EnQue<Y>(quantOutLocal);
488+ outQueue.FreeTensor(outLocal);
489+ } else {
490+ outQueue.EnQue<OUT_DTYPE>(outLocal);
491+ }
492+}
493+} // namespace AdaLayerNormNS
494+ 
495+#endif // ADA_LAYER_NORM_FULL_LOAD_H
Anorm/ada_layer_norm/op_kernel/arch35/ada_layer_norm_impl.h+316-0
@@ -0,0 +1,316 @@
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 ada_layer_norm_impl.h
13+ * \brief
14+ */
15+#ifndef ADA_LAYER_NORM_IMPL_H
16+#define ADA_LAYER_NORM_IMPL_H
17+ 
18+#include "ada_layer_norm_full_load.h"
19+#include "ada_layer_norm_welford.h"
20+ 
21+using namespace AdaLayerNormNS;
22+ 
23+template <typename T, typename U, typename Y, uint8_t OP_CODE>
24+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::ProcessWelfordUpdate(int64_t computeLength, int64_t welfordCount)
25+{
26+ LocalTensor<T> xLocal = xQueue.DeQue<T>();
27+ WelfordUpdateParam para;
28+ para.rnLength = 1;
29+ para.abLength = sliceSize;
30+ para.abComputeLength = computeLength;
31+ para.nRec = 1.0f / static_cast<float>(welfordCount);
32+ WelfordUpdate<T, float, false>(meanTmpLocal, varTmpLocal, meanTmpLocal, varTmpLocal, xLocal, tmpLocal, para);
33+ PipeBarrier<PIPE_V>();
34+ xQueue.FreeTensor(xLocal);
35+}
36+ 
37+template <typename T, typename U, typename Y, uint8_t OP_CODE>
38+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::ProcessWelfordFinalize(int64_t welfordCount, int64_t batchCount)
39+{
40+ WelfordFinalizePara para;
41+ para.rnLength = welfordCount;
42+ para.abLength = sliceSize;
43+ para.headCount = welfordCount;
44+ para.headCountLength = tailSize;
45+ para.tailCount = (tailSize > 0) ? (welfordCount - 1) : welfordCount;
46+ para.tailCountLength = sliceSize - tailSize;
47+ para.abRec = 1.0f / static_cast<float>(sliceSize);
48+ para.rRec = 1.0f / static_cast<float>(hiddenDim);
49+ WelfordFinalize<true>(meanLocal[batchCount], varLocal[batchCount], meanTmpLocal, varTmpLocal, tmpLocal, para);
50+ PipeBarrier<PIPE_V>();
51+ event_t eventIdVToMte2 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE2));
52+ SetFlag<HardEvent::V_MTE2>(eventIdVToMte2);
53+ WaitFlag<HardEvent::V_MTE2>(eventIdVToMte2);
54+}
55+ 
56+template <typename T, typename U, typename Y, uint8_t OP_CODE>
57+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::ProcessNormalize(int64_t dataCount, int64_t batchCount)
58+{
59+ LocalTensor<float> xLocal;
60+ if constexpr (std::is_same_v<T, float>) {
61+ xLocal = xQueue.DeQue<float>();
62+ } else {
63+ LocalTensor<T> xFp16 = xQueue.DeQue<T>();
64+ xLocal = xFp16.template ReinterpretCast<float>();
65+ Cast(xLocal, xFp16[DATA_COUNT], RoundMode::CAST_NONE, dataCount);
66+ PipeBarrier<PIPE_V>();
67+ }
68+ 
69+ NormalizePara para;
70+ para.aLength = 1;
71+ para.rLength = dataCount;
72+ para.rLengthWithPadding = DATA_COUNT;
73+ if (hasWeight && hasBias) {
74+ Normalize<float, float, false, hasGammaBetaNormalizeConfig>(
75+ normLocal, rstdLocal[batchCount], meanLocal[batchCount], varLocal[batchCount],
76+ xLocal, weightLocal, biasLocal, tmpLocal, epsilon, para);
77+ } else if (hasWeight && !hasBias) {
78+ Normalize<float, float, false, hasGammaNoBetaNormalizeConfig>(
79+ normLocal, rstdLocal[batchCount], meanLocal[batchCount], varLocal[batchCount],
80+ xLocal, weightLocal, biasLocal, tmpLocal, epsilon, para);
81+ } else if (!hasWeight && hasBias) {
82+ Normalize<float, float, false, noGammaHasBetaNormalizeConfig>(
83+ normLocal, rstdLocal[batchCount], meanLocal[batchCount], varLocal[batchCount],
84+ xLocal, weightLocal, biasLocal, tmpLocal, epsilon, para);
85+ } else if (!hasWeight && !hasBias) {
86+ Normalize<float, float, false, noGammaNoBetaNormalizeConfig>(
87+ normLocal, rstdLocal[batchCount], meanLocal[batchCount], varLocal[batchCount],
88+ xLocal, weightLocal, biasLocal, tmpLocal, epsilon, para);
89+ }
90+ PipeBarrier<PIPE_V>();
91+ event_t eventIdVToMte2 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::V_MTE2));
92+ SetFlag<HardEvent::V_MTE2>(eventIdVToMte2);
93+ WaitFlag<HardEvent::V_MTE2>(eventIdVToMte2);
94+ xQueue.FreeTensor(xLocal);
95+}
96+ 
97+template <typename T, typename U, typename Y, uint8_t OP_CODE>
98+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::CalculateScale(int64_t batchCount)
99+{
100+ __local_mem__ float* maxTmpAddr = (__ubuf__ float*)maxTmpLocal.GetPhyAddr();
101+ __local_mem__ float* quantScaleAddr = (__ubuf__ float*)quantScaleLocal[batchCount].GetPhyAddr();
102+ 
103+ __VEC_SCOPE__
104+ {
105+ RegTensor<float> quantScale;
106+ MaskReg pregMerge = CreateMask<float, MaskPattern::VL1>();
107+ 
108+ DataCopy<float, LoadDist::DIST_BRC_B32>(quantScale, maxTmpAddr);
109+ Muls(quantScale, quantScale, quantFactor, pregMerge);
110+ // 拷出量化系数
111+ DataCopy<float, StoreDist::DIST_FIRST_ELEMENT_B32>(quantScaleAddr, quantScale, pregMerge);
112+ }
113+}
114+ 
115+template <typename T, typename U, typename Y, uint8_t OP_CODE>
116+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::CopyInOtherData()
117+{
118+ if (hasWeight) {
119+ CopyInAndCast(weightLocal, weightGm, hiddenDim, DATA_COUNT);
120+ }
121+ if (hasBias) {
122+ CopyInAndCast(biasLocal, biasGm, hiddenDim, DATA_COUNT);
123+ }
124+ if constexpr (OP_CODE == QUANT_OP_CODE) {
125+ if (hasSmooth) {
126+ LocalTensor<T> smoothLocal = smoothQueue.AllocTensor<T>();
127+ DataCopyExtParams copyInParams{1, static_cast<uint32_t>(hiddenDim * sizeof(T)), 0, 0, 0};
128+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
129+ DataCopyPad(smoothLocal, smoothGm, copyInParams, padParams);
130+ smoothQueue.EnQue(smoothLocal);
131+ }
132+ }
133+}
134+ 
135+template <typename T, typename U, typename Y, uint8_t OP_CODE>
136+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::CopyInScaleShift(int64_t offset, uint16_t blockCount, int64_t len)
137+{
138+ LocalTensor<T> scaleLocal = scaleQueue.AllocTensor<T>();
139+ LocalTensor<T> shiftLocal = shiftQueue.AllocTensor<T>();
140+ DataCopyExtParams copyInParams{blockCount, static_cast<uint32_t>(len * sizeof(T)), 0, 0, 0};
141+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
142+ if (blockCount > 1 && len < hiddenDimCeil) {
143+ padParams.isPad = true;
144+ padParams.paddingValue = 0;
145+ padParams.rightPadding = hiddenDimCeil - len;
146+ }
147+ DataCopyPad(scaleLocal, scaleGm[offset], copyInParams, padParams);
148+ DataCopyPad(shiftLocal, shiftGm[offset], copyInParams, padParams);
149+ scaleQueue.EnQue(scaleLocal);
150+ shiftQueue.EnQue(shiftLocal);
151+}
152+ 
153+template <typename T, typename U, typename Y, uint8_t OP_CODE>
154+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::CopyInX(int64_t offset, uint16_t blockCount, int64_t len)
155+{
156+ LocalTensor<T> xLocal = xQueue.AllocTensor<T>();
157+ DataCopyExtParams copyInParams{blockCount, static_cast<uint32_t>(len * sizeof(T)), 0, 0, 0};
158+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
159+ if (blockCount > 1 && len < hiddenDimCeil) {
160+ padParams.isPad = true;
161+ padParams.paddingValue = 0;
162+ padParams.rightPadding = hiddenDimCeil - len;
163+ }
164+ if constexpr (std::is_same_v<T, float>) {
165+ DataCopyPad(xLocal, xGm[offset], copyInParams, padParams);
166+ } else {
167+ DataCopyPad(xLocal[DATA_COUNT], xGm[offset], copyInParams, padParams);
168+ }
169+ xQueue.EnQue(xLocal);
170+}
171+ 
172+template <typename T, typename U, typename Y, uint8_t OP_CODE>
173+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::BaseCopyOut(int64_t offset, uint16_t blockCount, int64_t len)
174+{
175+ LocalTensor<T> out = outQueue.DeQue<T>();
176+ DataCopyExtParams copyOutParams{blockCount, static_cast<uint32_t>(len * sizeof(T)), 0, 0, 0};
177+ DataCopyPad(outGm[offset], out, copyOutParams);
178+ outQueue.FreeTensor(out);
179+}
180+ 
181+template <typename T, typename U, typename Y, uint8_t OP_CODE>
182+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::QuantCopyOut(int64_t offset, uint16_t blockCount, int64_t len)
183+{
184+ LocalTensor<Y> quantOutLocal = quantOutQueue.DeQue<Y>();
185+ DataCopyExtParams copyOutParams{blockCount, static_cast<uint32_t>(len * sizeof(Y)), 0, 0, 0};
186+ DataCopyPad(quantOutGm[offset], quantOutLocal, copyOutParams);
187+ quantOutQueue.FreeTensor(quantOutLocal);
188+}
189+ 
190+template <typename T, typename U, typename Y, uint8_t OP_CODE>
191+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::CopyMeanRstdOut(int64_t offset, int64_t len)
192+{
193+ LocalTensor<T> meanOut = meanQueue.DeQue<T>();
194+ LocalTensor<T> rstdOut = rstdQueue.DeQue<T>();
195+ DataCopyExtParams copyOutParams{1, static_cast<uint32_t>(len * sizeof(T)), 0, 0, 0};
196+ DataCopyPad(meanGm[offset], meanOut, copyOutParams);
197+ DataCopyPad(rstdGm[offset], rstdOut, copyOutParams);
198+ meanQueue.FreeTensor(meanOut);
199+ rstdQueue.FreeTensor(rstdOut);
200+}
201+ 
202+template <typename T, typename U, typename Y, uint8_t OP_CODE>
203+__aicore__ inline void AdaLayerNormFullLoad<T, U, Y, OP_CODE>::CopyScaleOut(int64_t offset, int64_t len)
204+{
205+ LocalTensor<float> quantScaleLocal = quantScaleQueue.DeQue<float>();
206+ DataCopyExtParams copyOutParams{1, static_cast<uint32_t>(len * sizeof(float)), 0, 0, 0};
207+ DataCopyPad(quantScaleGm[offset], quantScaleLocal, copyOutParams);
208+ quantScaleQueue.FreeTensor(quantScaleLocal);
209+}
210+ 
211+template <typename T, typename U, typename Y, uint8_t OP_CODE>
212+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::CopyInOtherData(int64_t offset, int64_t len)
213+{
214+ if (hasWeight) {
215+ CopyInAndCast(weightLocal, weightGm[offset], len, DATA_COUNT);
216+ }
217+ if (hasBias) {
218+ CopyInAndCast(biasLocal, biasGm[offset], len, DATA_COUNT);
219+ }
220+ if constexpr (OP_CODE == QUANT_OP_CODE) {
221+ if (hasSmooth) {
222+ LocalTensor<T> smoothLocal = smoothQueue.AllocTensor<T>();
223+ DataCopyExtParams copyInParams{1, static_cast<uint32_t>(len * sizeof(T)), 0, 0, 0};
224+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
225+ DataCopyPad(smoothLocal, smoothGm[offset], copyInParams, padParams);
226+ smoothQueue.EnQue(smoothLocal);
227+ }
228+ }
229+}
230+ 
231+template <typename T, typename U, typename Y, uint8_t OP_CODE>
232+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::CopyInScaleShift(int64_t offset, int64_t len)
233+{
234+ LocalTensor<T> scaleLocal = scaleQueue.AllocTensor<T>();
235+ LocalTensor<T> shiftLocal = shiftQueue.AllocTensor<T>();
236+ DataCopyExtParams copyInParams{1, static_cast<uint32_t>(len * sizeof(T)), 0, 0, 0};
237+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
238+ DataCopyPad(scaleLocal, scaleGm[offset], copyInParams, padParams);
239+ DataCopyPad(shiftLocal, shiftGm[offset], copyInParams, padParams);
240+ scaleQueue.EnQue(scaleLocal);
241+ shiftQueue.EnQue(shiftLocal);
242+}
243+ 
244+template <typename T, typename U, typename Y, uint8_t OP_CODE>
245+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::CopyInX(int64_t offset, int64_t len)
246+{
247+ LocalTensor<T> xLocal = xQueue.AllocTensor<T>();
248+ DataCopyExtParams copyInParams{1, static_cast<uint32_t>(len * sizeof(T)), 0, 0, 0};
249+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
250+ if constexpr (std::is_same_v<T, float>) {
251+ DataCopyPad(xLocal, xGm[offset], copyInParams, padParams);
252+ } else {
253+ DataCopyPad(xLocal[DATA_COUNT], xGm[offset], copyInParams, padParams);
254+ }
255+ xQueue.EnQue(xLocal);
256+}
257+ 
258+template <typename T, typename U, typename Y, uint8_t OP_CODE>
259+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::CopyInNorm(int64_t offset, int64_t len)
260+{
261+ LocalTensor<float> xLocal = xQueue.AllocTensor<float>();
262+ DataCopyExtParams copyInParams{1, static_cast<uint32_t>(len * sizeof(float)), 0, 0, 0};
263+ DataCopyPadExtParams<float> padParams{false, 0, 0, 0};
264+ DataCopyPad(xLocal, normGm[offset], copyInParams, padParams);
265+ xQueue.EnQue(xLocal);
266+}
267+ 
268+template <typename T, typename U, typename Y, uint8_t OP_CODE>
269+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::BaseCopyOut(int64_t offset, int64_t len)
270+{
271+ LocalTensor<T> out = outQueue.DeQue<T>();
272+ DataCopyExtParams copyOutParams{1, static_cast<uint32_t>(len * sizeof(T)), 0, 0, 0};
273+ DataCopyPad(outGm[offset], out, copyOutParams);
274+ outQueue.FreeTensor(out);
275+}
276+ 
277+template <typename T, typename U, typename Y, uint8_t OP_CODE>
278+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::CopyNormOut(int64_t offset, int64_t len)
279+{
280+ LocalTensor<float> out = outQueue.DeQue<float>();
281+ DataCopyExtParams copyOutParams{1, static_cast<uint32_t>(len * sizeof(float)), 0, 0, 0};
282+ DataCopyPad(normGm[offset], out, copyOutParams);
283+ outQueue.FreeTensor(out);
284+}
285+ 
286+template <typename T, typename U, typename Y, uint8_t OP_CODE>
287+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::QuantCopyOut(int64_t offset, int64_t len)
288+{
289+ LocalTensor<Y> quantOutLocal = quantOutQueue.DeQue<Y>();
290+ DataCopyExtParams copyOutParams{1, static_cast<uint32_t>(len * sizeof(Y)), 0, 0, 0};
291+ DataCopyPad(quantOutGm[offset], quantOutLocal, copyOutParams);
292+ quantOutQueue.FreeTensor(quantOutLocal);
293+}
294+ 
295+template <typename T, typename U, typename Y, uint8_t OP_CODE>
296+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::CopyScaleOut(int64_t offset, int64_t len)
297+{
298+ LocalTensor<float> quantScaleLocal = quantScaleQueue.DeQue<float>();
299+ DataCopyExtParams copyOutParams{1, static_cast<uint32_t>(len * sizeof(float)), 0, 0, 0};
300+ DataCopyPad(quantScaleGm[offset], quantScaleLocal, copyOutParams);
301+ quantScaleQueue.FreeTensor(quantScaleLocal);
302+}
303+ 
304+template <typename T, typename U, typename Y, uint8_t OP_CODE>
305+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::CopyMeanRstdOut(int64_t offset, int64_t len)
306+{
307+ LocalTensor<T> meanOut = meanQueue.DeQue<T>();
308+ LocalTensor<T> rstdOut = rstdQueue.DeQue<T>();
309+ DataCopyExtParams copyOutParams{1, static_cast<uint32_t>(len * sizeof(T)), 0, 0, 0};
310+ DataCopyPad(meanGm[offset], meanOut, copyOutParams);
311+ DataCopyPad(rstdGm[offset], rstdOut, copyOutParams);
312+ meanQueue.FreeTensor(meanOut);
313+ rstdQueue.FreeTensor(rstdOut);
314+}
315+ 
316+#endif // ADA_LAYER_NORM_IMPL_H
Anorm/ada_layer_norm/op_kernel/arch35/ada_layer_norm_welford.h+543-0
@@ -0,0 +1,543 @@
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 ada_layer_norm_welford.h
13+ * \brief
14+ */
15+#ifndef ADA_LAYER_NORM_WELFORD_H
16+#define ADA_LAYER_NORM_WELFORD_H
17+ 
18+#include "ada_layer_norm_common.h"
19+ 
20+namespace AdaLayerNormNS {
21+using namespace AscendC;
22+ 
23+template <typename T, typename U, typename Y, uint8_t OP_CODE>
24+class AdaLayerNormWelford {
25+public:
26+ __aicore__ inline AdaLayerNormWelford(){};
27+ __aicore__ inline void InitV2(const GmAddr* gmAddr, const AdaLayerNormTilingData *tilingData);
28+ __aicore__ inline void InitQuant(const GmAddr* gmAddr, GM_ADDR workspace, const AdaLayerNormTilingData *tilingData);
29+ __aicore__ inline void Process();
30+ 
31+private:
32+ __aicore__ inline void ParseTilingData(const AdaLayerNormTilingData *tilingData);
33+ __aicore__ inline void SliceProcess();
34+ __aicore__ inline void ComputeMeanVar(int64_t offset, int64_t batchCount);
35+ __aicore__ inline void ComputeAdaLayerNorm(int64_t offset, int64_t scaleOffset, int64_t batchCount);
36+ __aicore__ inline void ProcessWelfordUpdate(int64_t computeLength, int64_t welfordCount);
37+ __aicore__ inline void ProcessWelfordFinalize(int64_t welfordCount, int64_t batchCount);
38+ __aicore__ inline void ProcessNormalize(int64_t dataCount, int64_t batchCount);
39+ __aicore__ inline void Adaption(int64_t dataCount, int64_t batchCount);
40+ __aicore__ inline void DynamicQuant(int64_t offset, int64_t batchCount);
41+ __aicore__ inline void CalculateScale(int64_t batchCount);
42+ __aicore__ inline void ProcessQuant(int64_t dataCount, int64_t batchCount);
43+ 
44+ __aicore__ inline void CopyInOtherData(int64_t offset, int64_t len);
45+ __aicore__ inline void CopyInScaleShift(int64_t offset, int64_t len);
46+ __aicore__ inline void CopyInX(int64_t offset, int64_t len);
47+ __aicore__ inline void CopyInNorm(int64_t offset, int64_t len);
48+ __aicore__ inline void BaseCopyOut(int64_t offset, int64_t len);
49+ __aicore__ inline void CopyNormOut(int64_t offset, int64_t len);
50+ __aicore__ inline void QuantCopyOut(int64_t offset, int64_t len);
51+ __aicore__ inline void CopyMeanRstdOut(int64_t offset, int64_t len);
52+ __aicore__ inline void CopyScaleOut(int64_t offset, int64_t len);
53+ 
54+private:
55+ using OUT_DTYPE = std::conditional_t<OP_CODE == QUANT_OP_CODE, float, T>;
56+ TPipe pipe;
57+ TQue<QuePosition::VECIN, 1> xQueue;
58+ TQue<QuePosition::VECIN, 1> scaleQueue;
59+ TQue<QuePosition::VECIN, 1> shiftQueue;
60+ TQue<QuePosition::VECIN, 1> smoothQueue;
61+ TBuf<TPosition::VECCALC> weightBuf;
62+ TBuf<TPosition::VECCALC> biasBuf;
63+ TBuf<TPosition::VECCALC> normBuf;
64+ TBuf<TPosition::VECCALC> tmpBuf;
65+ TBuf<TPosition::VECCALC> varBuf;
66+ TBuf<TPosition::VECCALC> maxBuf;
67+ TQue<QuePosition::VECOUT, 1> outQueue;
68+ TQue<QuePosition::VECOUT, 1> quantOutQueue;
69+ TQue<QuePosition::VECOUT, 1> meanQueue;
70+ TQue<QuePosition::VECOUT, 1> rstdQueue;
71+ TQue<QuePosition::VECOUT, 1> quantScaleQueue;
72+ 
73+ GlobalTensor<T> xGm;
74+ GlobalTensor<T> scaleGm;
75+ GlobalTensor<T> shiftGm;
76+ GlobalTensor<U> weightGm;
77+ GlobalTensor<U> biasGm;
78+ GlobalTensor<T> smoothGm;
79+ GlobalTensor<T> outGm;
80+ GlobalTensor<T> meanGm;
81+ GlobalTensor<T> rstdGm;
82+ GlobalTensor<Y> quantOutGm;
83+ GlobalTensor<float> quantScaleGm;
84+ GlobalTensor<float> normGm;
85+ 
86+ LocalTensor<float> weightLocal;
87+ LocalTensor<float> biasLocal;
88+ LocalTensor<float> normLocal;
89+ LocalTensor<uint8_t> tmpLocal;
90+ LocalTensor<float> meanTmpLocal;
91+ LocalTensor<float> varTmpLocal;
92+ LocalTensor<float> varLocal;
93+ LocalTensor<float> maxTmpLocal;
94+ LocalTensor<float> meanLocal;
95+ LocalTensor<float> rstdLocal;
96+ LocalTensor<float> quantScaleLocal;
97+ 
98+ int64_t sliceSize = 0;
99+ int64_t sliceCount = 0;
100+ int64_t tailSize = 0;
101+ int64_t seqLen = 0;
102+ int64_t hiddenDim = 0;
103+ int64_t hiddenDimCeil = 0;
104+ int64_t outIntCeil = 0;
105+ int64_t rowStart = 0;
106+ int64_t rowEnd = 0;
107+ int64_t normOffset = 0;
108+ int64_t tmpBufferSize = 0;
109+ float epsilon = 0.0f;
110+ int32_t hasWeight = 0;
111+ int32_t hasBias = 0;
112+ int32_t hasSmooth = 0;
113+ float quantFactor = 0.0f;
114+};
115+ 
116+template <typename T, typename U, typename Y, uint8_t OP_CODE>
117+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::InitV2(const GmAddr* gmAddr, const AdaLayerNormTilingData *tilingData)
118+{
119+ ParseTilingData(tilingData);
120+ pipe.InitBuffer(xQueue, 1, DATA_COUNT * sizeof(float));
121+ pipe.InitBuffer(scaleQueue, 1, DATA_COUNT * sizeof(float));
122+ pipe.InitBuffer(shiftQueue, 1, DATA_COUNT * sizeof(float));
123+ pipe.InitBuffer(weightBuf, DATA_COUNT * sizeof(float));
124+ pipe.InitBuffer(biasBuf, DATA_COUNT * sizeof(float));
125+ pipe.InitBuffer(varBuf, BATCH_COUNT * sizeof(float));
126+ pipe.InitBuffer(normBuf, DATA_COUNT * sizeof(float));
127+ if (tmpBufferSize > 0) {
128+ pipe.InitBuffer(tmpBuf, tmpBufferSize);
129+ }
130+ pipe.InitBuffer(outQueue, 1, DATA_COUNT * sizeof(float));
131+ pipe.InitBuffer(meanQueue, 1, BATCH_COUNT * sizeof(float));
132+ pipe.InitBuffer(rstdQueue, 1, BATCH_COUNT * sizeof(float));
133+ 
134+ xGm.SetGlobalBuffer((__gm__ T*)gmAddr->x);
135+ scaleGm.SetGlobalBuffer((__gm__ T*)gmAddr->scale);
136+ shiftGm.SetGlobalBuffer((__gm__ T*)gmAddr->shift);
137+ weightGm.SetGlobalBuffer((__gm__ U*)gmAddr->weight);
138+ biasGm.SetGlobalBuffer((__gm__ U*)gmAddr->bias);
139+ outGm.SetGlobalBuffer((__gm__ T*)gmAddr->out);
140+ meanGm.SetGlobalBuffer((__gm__ T*)gmAddr->mean);
141+ rstdGm.SetGlobalBuffer((__gm__ T*)gmAddr->rstd);
142+}
143+ 
144+template <typename T, typename U, typename Y, uint8_t OP_CODE>
145+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::InitQuant(const GmAddr* gmAddr, GM_ADDR workspace,
146+ const AdaLayerNormTilingData *tilingData)
147+{
148+ ParseTilingData(tilingData);
149+ pipe.InitBuffer(xQueue, 1, DATA_COUNT * sizeof(float));
150+ pipe.InitBuffer(scaleQueue, 1, DATA_COUNT * sizeof(float));
151+ pipe.InitBuffer(shiftQueue, 1, DATA_COUNT * sizeof(float));
152+ pipe.InitBuffer(smoothQueue, 1, DATA_COUNT * sizeof(float));
153+ pipe.InitBuffer(weightBuf, DATA_COUNT * sizeof(float));
154+ pipe.InitBuffer(biasBuf, DATA_COUNT * sizeof(float));
155+ pipe.InitBuffer(varBuf, BATCH_COUNT * sizeof(float));
156+ pipe.InitBuffer(maxBuf, V_LENGTH * sizeof(float));
157+ pipe.InitBuffer(normBuf, DATA_COUNT * sizeof(float));
158+ if (tmpBufferSize > 0) {
159+ pipe.InitBuffer(tmpBuf, tmpBufferSize);
160+ }
161+ pipe.InitBuffer(outQueue, 1, DATA_COUNT * sizeof(float));
162+ pipe.InitBuffer(quantOutQueue, 1, DATA_COUNT * sizeof(int8_t));
163+ pipe.InitBuffer(meanQueue, 1, BATCH_COUNT * sizeof(float));
164+ pipe.InitBuffer(rstdQueue, 1, BATCH_COUNT * sizeof(float));
165+ pipe.InitBuffer(quantScaleQueue, 1, BATCH_COUNT * sizeof(float));
166+ 
167+ xGm.SetGlobalBuffer((__gm__ T*)gmAddr->x);
168+ scaleGm.SetGlobalBuffer((__gm__ T*)gmAddr->scale);
169+ shiftGm.SetGlobalBuffer((__gm__ T*)gmAddr->shift);
170+ weightGm.SetGlobalBuffer((__gm__ U*)gmAddr->weight);
171+ biasGm.SetGlobalBuffer((__gm__ U*)gmAddr->bias);
172+ smoothGm.SetGlobalBuffer((__gm__ T*)gmAddr->smooth_scales);
173+ normGm.SetGlobalBuffer((__gm__ float*)workspace);
174+ quantOutGm.SetGlobalBuffer((__gm__ Y*)gmAddr->out);
175+ quantScaleGm.SetGlobalBuffer((__gm__ float*)gmAddr->quant_scale);
176+}
177+ 
178+template <typename T, typename U, typename Y, uint8_t OP_CODE>
179+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::Process()
180+{
181+ if (rowStart >= rowEnd) {
182+ return;
183+ }
184+ 
185+ weightLocal = weightBuf.Get<float>();
186+ biasLocal = biasBuf.Get<float>();
187+ normLocal = normBuf.Get<float>();
188+ tmpLocal = tmpBuf.Get<uint8_t>();
189+ meanTmpLocal = weightLocal;
190+ varTmpLocal = biasLocal;
191+ varLocal = varBuf.Get<float>();
192+ meanLocal = meanQueue.AllocTensor<float>();
193+ rstdLocal = rstdQueue.AllocTensor<float>();
194+ if constexpr (OP_CODE == QUANT_OP_CODE) {
195+ quantScaleLocal = quantScaleQueue.AllocTensor<float>();
196+ maxTmpLocal = maxBuf.Get<float>();
197+ }
198+ SliceProcess();
199+}
200+ 
201+template <typename T, typename U, typename Y, uint8_t OP_CODE>
202+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::ParseTilingData(const AdaLayerNormTilingData *tilingData)
203+{
204+ int32_t blockNum = HALF_BLOCK_NUM;
205+ if constexpr (std::is_same_v<T, float>) {
206+ blockNum = FLOAT_BLOCK_NUM;
207+ }
208+ quantFactor = GetQuantFactor<Y>();
209+ seqLen = tilingData->seqLen;
210+ hiddenDim = tilingData->hiddenDim;
211+ hiddenDimCeil = CeilA2B(hiddenDim, blockNum) * blockNum;
212+ outIntCeil = CeilA2B(hiddenDim, INT8_BLOCK_NUM) * INT8_BLOCK_NUM;
213+ epsilon = tilingData->epsilon;
214+ hasWeight = tilingData->hasWeight;
215+ hasBias = tilingData->hasBias;
216+ hasSmooth = tilingData->hasSmooth;
217+ sliceSize = tilingData->sliceSize;
218+ sliceCount = hiddenDim / sliceSize;
219+ tailSize = hiddenDim % sliceSize;
220+ tmpBufferSize = tilingData->tmpBufferSize;
221+ 
222+ int64_t singleCoreNum = tilingData->singleCoreNum;
223+ int64_t tailNum = tilingData->tailNum;
224+ int64_t blockIdx = GetBlockIdx();
225+ rowStart = singleCoreNum * blockIdx + (blockIdx < tailNum ? blockIdx : tailNum);
226+ rowEnd = rowStart + singleCoreNum + (blockIdx < tailNum ? 1 : 0);
227+ normOffset = blockIdx * outIntCeil;
228+}
229+ 
230+template <typename T, typename U, typename Y, uint8_t OP_CODE>
231+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::SliceProcess()
232+{
233+ int64_t batchCount = 0;
234+ int64_t rowIdx = rowStart;
235+ while (rowIdx < rowEnd) {
236+ int64_t offset = rowIdx * hiddenDim;
237+ int64_t scaleOffset = (rowIdx / seqLen) * hiddenDim;
238+ // 计算均值和方差
239+ ComputeMeanVar(offset, batchCount);
240+ ComputeAdaLayerNorm(offset, scaleOffset, batchCount);
241+ if constexpr (OP_CODE == QUANT_OP_CODE) {
242+ DynamicQuant(offset, batchCount);
243+ }
244+ batchCount++;
245+ rowIdx++;
246+ if (batchCount == BATCH_COUNT || rowIdx == rowEnd) {
247+ if constexpr (OP_CODE == QUANT_OP_CODE) {
248+ quantScaleQueue.EnQue<float>(quantScaleLocal);
249+ CopyScaleOut(rowIdx - batchCount, batchCount);
250+ if (rowIdx != rowEnd) {
251+ quantScaleLocal = quantScaleQueue.AllocTensor<float>();
252+ }
253+ } else {
254+ if constexpr (!std::is_same_v<T, float>) {
255+ LocalTensor<T> meanOut = meanLocal.ReinterpretCast<T>();
256+ LocalTensor<T> rstdOut = rstdLocal.ReinterpretCast<T>();
257+ Cast(meanOut, meanLocal, RoundMode::CAST_RINT, batchCount);
258+ Cast(rstdOut, rstdLocal, RoundMode::CAST_RINT, batchCount);
259+ meanQueue.EnQue<T>(meanOut);
260+ rstdQueue.EnQue<T>(rstdOut);
261+ } else {
262+ meanQueue.EnQue<float>(meanLocal);
263+ rstdQueue.EnQue<float>(rstdLocal);
264+ }
265+ CopyMeanRstdOut(rowIdx - batchCount, batchCount);
266+ if (rowIdx < rowEnd) {
267+ meanLocal = meanQueue.AllocTensor<float>();
268+ rstdLocal = rstdQueue.AllocTensor<float>();
269+ }
270+ }
271+ batchCount = 0;
272+ }
273+ }
274+}
275+ 
276+template <typename T, typename U, typename Y, uint8_t OP_CODE>
277+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::ComputeMeanVar(int64_t offset, int64_t batchCount)
278+{
279+ Duplicate(meanTmpLocal, 0.0f, sliceSize);
280+ Duplicate(varTmpLocal, 0.0f, sliceSize);
281+ PipeBarrier<PIPE_V>();
282+ int64_t welfordCount = 0;
283+ while (welfordCount < sliceCount) {
284+ welfordCount ++;
285+ LocalTensor<T> xLocal = xQueue.AllocTensor<T>();
286+ DataCopyPad(xLocal, xGm[offset], {1, static_cast<uint32_t>(sliceSize * sizeof(T)), 0, 0, 0}, {false, 0, 0, 0});
287+ xQueue.EnQue(xLocal);
288+ ProcessWelfordUpdate(sliceSize, welfordCount);
289+ offset += sliceSize;
290+ }
291+ if (tailSize > 0) {
292+ welfordCount ++;
293+ LocalTensor<T> xLocal = xQueue.AllocTensor<T>();
294+ DataCopyPad(xLocal, xGm[offset], {1, static_cast<uint32_t>(tailSize * sizeof(T)), 0, 0, 0}, {false, 0, 0, 0});
295+ xQueue.EnQue(xLocal);
296+ ProcessWelfordUpdate(tailSize, welfordCount);
297+ }
298+ ProcessWelfordFinalize(welfordCount, batchCount);
299+}
300+ 
301+template <typename T, typename U, typename Y, uint8_t OP_CODE>
302+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::ComputeAdaLayerNorm(int64_t offset, int64_t scaleOffset, int64_t batchCount)
303+{
304+ if constexpr (OP_CODE == QUANT_OP_CODE) {
305+ Duplicate(maxTmpLocal, 0.0f, sliceSize);
306+ PipeBarrier<PIPE_V>();
307+ }
308+ int64_t h = 0;
309+ for (int64_t i = 0; i < sliceCount; i ++) {
310+ CopyInOtherData(h, sliceSize);
311+ CopyInX(offset + h, sliceSize);
312+ ProcessNormalize(sliceSize, batchCount);
313+ CopyInScaleShift(scaleOffset + h, sliceSize);
314+ Adaption(sliceSize, batchCount);
315+ if constexpr (OP_CODE == QUANT_OP_CODE) {
316+ CopyNormOut(normOffset + h, sliceSize);
317+ } else {
318+ BaseCopyOut(offset + h, sliceSize);
319+ }
320+ h += sliceSize;
321+ }
322+ if (tailSize > 0) {
323+ CopyInOtherData(h, tailSize);
324+ CopyInX(offset + h, tailSize);
325+ ProcessNormalize(tailSize, batchCount);
326+ CopyInScaleShift(scaleOffset + h, tailSize);
327+ Adaption(tailSize, batchCount);
328+ if constexpr (OP_CODE == QUANT_OP_CODE) {
329+ CopyNormOut(normOffset + h, tailSize);
330+ } else {
331+ BaseCopyOut(offset + h, tailSize);
332+ }
333+ }
334+}
335+ 
336+template <typename T, typename U, typename Y, uint8_t OP_CODE>
337+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::DynamicQuant(int64_t offset, int64_t batchCount)
338+{
339+ CalculateScale(batchCount);
340+ event_t eventIdMte3ToMte2 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE3_MTE2));
341+ SetFlag<HardEvent::MTE3_MTE2>(eventIdMte3ToMte2);
342+ WaitFlag<HardEvent::MTE3_MTE2>(eventIdMte3ToMte2);
343+ int64_t h = 0;
344+ for (int64_t i = 0; i < sliceCount; i ++) {
345+ CopyInNorm(normOffset + h, sliceSize);
346+ ProcessQuant(sliceSize, batchCount);
347+ QuantCopyOut(offset + h, sliceSize);
348+ h += sliceSize;
349+ }
350+ if (tailSize > 0) {
351+ CopyInNorm(normOffset + h, tailSize);
352+ ProcessQuant(tailSize, batchCount);
353+ QuantCopyOut(offset + h, tailSize);
354+ }
355+}
356+ 
357+template <typename T, typename U, typename Y, uint8_t OP_CODE>
358+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::Adaption(int64_t dataCount, int64_t batchCount)
359+{
360+ LocalTensor<T> scaleLocal = scaleQueue.DeQue<T>();
361+ LocalTensor<T> shiftLocal = shiftQueue.DeQue<T>();
362+ LocalTensor<T> smoothLocal;
363+ LocalTensor<OUT_DTYPE> outLocal = outQueue.AllocTensor<OUT_DTYPE>();
364+ 
365+ __local_mem__ float* normAddr = (__ubuf__ float*)normLocal.GetPhyAddr();
366+ __local_mem__ OUT_DTYPE* outAddr = (__ubuf__ OUT_DTYPE*)outLocal.GetPhyAddr();
367+ __local_mem__ T* scaleAddr = (__ubuf__ T*)scaleLocal.GetPhyAddr();
368+ __local_mem__ T* shiftAddr = (__ubuf__ T*)shiftLocal.GetPhyAddr();
369+ 
370+ __local_mem__ T* smoothAddr;
371+ __local_mem__ float* maxTmpAddr;
372+ if constexpr (OP_CODE == QUANT_OP_CODE) {
373+ if (hasSmooth) {
374+ smoothLocal = smoothQueue.DeQue<T>();
375+ smoothAddr = (__ubuf__ T*)smoothLocal.GetPhyAddr();
376+ }
377+ maxTmpAddr = (__ubuf__ float*)maxTmpLocal.GetPhyAddr();
378+ }
379+ 
380+ uint16_t tailLength = static_cast<uint16_t>(dataCount % TWO_V_LENGTH);
381+ uint16_t colLoopTimes = static_cast<uint16_t>(dataCount / TWO_V_LENGTH) + (tailLength > V_LENGTH ? 1 : 0);
382+ uint32_t rightLength = static_cast<uint32_t>(dataCount - colLoopTimes * V_LENGTH);
383+ __VEC_SCOPE__
384+ {
385+ RegTensor<float> x1;
386+ RegTensor<float> x2;
387+ RegTensor<float> scale1;
388+ RegTensor<float> scale2;
389+ RegTensor<float> shift1;
390+ RegTensor<float> shift2;
391+ RegTensor<float> tmpMax;
392+ 
393+ MaskReg pregFull = CreateMask<float, MaskPattern::ALL>();
394+ MaskReg pregMerge = CreateMask<float, MaskPattern::VL1>();
395+ MaskReg pregLoop;
396+ uint32_t sreg2 = rightLength;
397+ if constexpr (OP_CODE == QUANT_OP_CODE) {
398+ DataCopy<float, LoadDist::DIST_BRC_B32>(tmpMax, maxTmpAddr);
399+ }
400+ for (uint16_t j = 0; j < colLoopTimes;j ++) {
401+ pregLoop = UpdateMask<float>(sreg2);
402+ DataCopy(x1, normAddr + j * TWO_V_LENGTH);
403+ DataCopy(x2, normAddr + j * TWO_V_LENGTH + V_LENGTH);
404+ LoadTensor(scale1, scaleAddr + j * TWO_V_LENGTH, pregFull);
405+ LoadTensor(scale2, scaleAddr + j * TWO_V_LENGTH + V_LENGTH, pregLoop);
406+ LoadTensor(shift1, shiftAddr + j * TWO_V_LENGTH, pregFull);
407+ LoadTensor(shift2, shiftAddr + j * TWO_V_LENGTH + V_LENGTH, pregLoop);
408+ Adds(scale1, scale1, 1.0f, pregFull);
409+ Adds(scale2, scale2, 1.0f, pregLoop);
410+ FusedMulDstAdd(x1, scale1, shift1, pregFull);
411+ FusedMulDstAdd(x2, scale2, shift2, pregLoop);
412+ if constexpr (OP_CODE == QUANT_OP_CODE) {
413+ if (hasSmooth) {
414+ RegTensor<float> smooth1;
415+ RegTensor<float> smooth2;
416+ LoadTensor(smooth1, smoothAddr + j * TWO_V_LENGTH, pregFull);
417+ LoadTensor(smooth2, smoothAddr + j * TWO_V_LENGTH + V_LENGTH, pregLoop);
418+ Mul(x1, x1, smooth1, pregFull);
419+ Mul(x2, x2, smooth2, pregLoop);
420+ }
421+ CopyToTensor(outAddr + j * TWO_V_LENGTH, x1, pregFull);
422+ CopyToTensor(outAddr + j * TWO_V_LENGTH + V_LENGTH, x2, pregLoop);
423+ Abs(x1, x1, pregFull);
424+ Abs(x2, x2, pregLoop);
425+ Max(x1, x1, x2, pregFull);
426+ Max(tmpMax, tmpMax, x1, pregFull);
427+ } else {
428+ CopyToTensor(outAddr + j * TWO_V_LENGTH, x1, pregFull);
429+ CopyToTensor(outAddr + j * TWO_V_LENGTH + V_LENGTH, x2, pregLoop);
430+ }
431+ }
432+ if (tailLength > 0 && tailLength <= V_LENGTH) {
433+ pregLoop = UpdateMask<float>(sreg2);
434+ DataCopy(x1, normAddr + colLoopTimes * TWO_V_LENGTH);
435+ LoadTensor(scale1, scaleAddr + colLoopTimes * TWO_V_LENGTH, pregLoop);
436+ LoadTensor(shift1, shiftAddr + colLoopTimes * TWO_V_LENGTH, pregLoop);
437+ Adds(scale1, scale1, 1.0f, pregLoop);
438+ FusedMulDstAdd(x1, scale1, shift1, pregLoop);
439+ if constexpr (OP_CODE == QUANT_OP_CODE) {
440+ if (hasSmooth) {
441+ RegTensor<float> smooth1;
442+ LoadTensor(smooth1, smoothAddr + colLoopTimes * TWO_V_LENGTH, pregLoop);
443+ Mul(x1, x1, smooth1, pregLoop);
444+ }
445+ CopyToTensor(outAddr + colLoopTimes * TWO_V_LENGTH, x1, pregLoop);
446+ Abs(x1, x1, pregLoop);
447+ Max(tmpMax, tmpMax, x1, pregFull);
448+ } else {
449+ CopyToTensor(outAddr + colLoopTimes * TWO_V_LENGTH, x1, pregLoop);
450+ }
451+ }
452+ if constexpr (OP_CODE == QUANT_OP_CODE) {
453+ ReduceMax(tmpMax, tmpMax, pregFull);
454+ DataCopy<float, StoreDist::DIST_FIRST_ELEMENT_B32>(maxTmpAddr, tmpMax, pregMerge);
455+ }
456+ }
457+ scaleQueue.FreeTensor(scaleLocal);
458+ shiftQueue.FreeTensor(shiftLocal);
459+ if constexpr (OP_CODE == QUANT_OP_CODE) {
460+ if (hasSmooth) {
461+ smoothQueue.FreeTensor(smoothLocal);
462+ }
463+ }
464+ outQueue.EnQue<OUT_DTYPE>(outLocal);
465+}
466+ 
467+template <typename T, typename U, typename Y, uint8_t OP_CODE>
468+__aicore__ inline void AdaLayerNormWelford<T, U, Y, OP_CODE>::ProcessQuant(int64_t dataCount, int64_t batchCount)
469+{
470+ LocalTensor<float> xLocal = xQueue.DeQue<float>();
471+ LocalTensor<Y> quantOutLocal = quantOutQueue.AllocTensor<Y>();
472+ 
473+ __local_mem__ float* xAddr = (__ubuf__ float*)xLocal.GetPhyAddr();
474+ __local_mem__ Y* quantOutAddr = (__ubuf__ Y*)quantOutLocal.GetPhyAddr();
475+ __local_mem__ float* quantScaleAddr = (__ubuf__ float*)quantScaleLocal[batchCount].GetPhyAddr();
476+ 
477+ uint16_t tailLength = static_cast<uint16_t>(dataCount % TWO_V_LENGTH);
478+ uint16_t colLoopTimes = static_cast<uint16_t>(dataCount / TWO_V_LENGTH) + (tailLength > V_LENGTH ? 1 : 0);
479+ uint32_t rightLength = static_cast<uint32_t>(dataCount - colLoopTimes * V_LENGTH);
480+ __VEC_SCOPE__
481+ {
482+ RegTensor<float> x1;
483+ RegTensor<float> x2;
484+ RegTensor<float> quantScale;
485+ RegTensor<int16_t> y1Int16;
486+ RegTensor<int16_t> y2Int16;
487+ RegTensor<half> y1Fp16;
488+ RegTensor<half> y2Fp16;
489+ RegTensor<Y> y1;
490+ RegTensor<Y> y2;
491+ 
492+ MaskReg pregFull = CreateMask<float, MaskPattern::ALL>();
493+ MaskReg pregMerge = CreateMask<float, MaskPattern::VL1>();
494+ MaskReg pregLoop;
495+ 
496+ DataCopy<float, LoadDist::DIST_BRC_B32>(quantScale, quantScaleAddr);
497+ uint32_t sreg2 = rightLength;
498+ for (uint16_t j = 0; j < colLoopTimes;j ++) {
499+ pregLoop = UpdateMask<float>(sreg2);
500+ DataCopy(x1, xAddr + j * TWO_V_LENGTH);
501+ DataCopy(x2, xAddr + j * TWO_V_LENGTH + V_LENGTH);
502+ Div(x1, x1, quantScale, pregFull);
503+ Div(x2, x2, quantScale, pregLoop);
504+ if constexpr (std::is_same_v<Y, int8_t>) {
505+ Cast<int16_t, float, castTraitF32ToI16>(y1Int16, x1, pregFull);
506+ Cast<int16_t, float, castTraitF32ToI16>(y2Int16, x2, pregLoop);
507+ Cast<half, int16_t, castTraitI16ToF16>(y1Fp16, y1Int16, pregFull);
508+ Cast<half, int16_t, castTraitI16ToF16>(y2Fp16, y2Int16, pregLoop);
509+ Cast<Y, half, castTraitF16ToI8>(y1, y1Fp16, pregFull);
510+ Cast<Y, half, castTraitF16ToI8>(y2, y2Fp16, pregLoop);
511+ } else if constexpr (std::is_same_v<Y, hifloat8_t>) {
512+ Cast<Y, float, castTraitF32Toh8>(y1, x1, pregFull);
513+ Cast<Y, float, castTraitF32Toh8>(y2, x2, pregLoop);
514+ } else {
515+ Cast<Y, float, castTraitF32Tofp8>(y1, x1, pregFull);
516+ Cast<Y, float, castTraitF32Tofp8>(y2, x2, pregLoop);
517+ }
518+ DataCopy<Y, StoreDist::DIST_PACK4_B32>(quantOutAddr + j * TWO_V_LENGTH, y1, pregFull);
519+ DataCopy<Y, StoreDist::DIST_PACK4_B32>(quantOutAddr + j * TWO_V_LENGTH + V_LENGTH, y2, pregLoop);
520+ }
521+ if (tailLength > 0 && tailLength <= V_LENGTH) {
522+ pregLoop = UpdateMask<float>(sreg2);
523+ DataCopy(x1, xAddr + colLoopTimes * TWO_V_LENGTH);
524+ Div(x1, x1, quantScale, pregLoop);
525+ if constexpr (std::is_same_v<Y, int8_t>) {
526+ Cast<int16_t, float, castTraitF32ToI16>(y1Int16, x1, pregLoop);
527+ Cast<half, int16_t, castTraitI16ToF16>(y1Fp16, y1Int16, pregLoop);
528+ Cast<Y, half, castTraitF16ToI8>(y1, y1Fp16, pregLoop);
529+ } else if constexpr (std::is_same_v<Y, hifloat8_t>) {
530+ Cast<Y, float, castTraitF32Toh8>(y1, x1, pregLoop);
531+ } else {
532+ Cast<Y, float, castTraitF32Tofp8>(y1, x1, pregLoop);
533+ }
534+ DataCopy<Y, StoreDist::DIST_PACK4_B32>(quantOutAddr + colLoopTimes * TWO_V_LENGTH, y1, pregLoop);
535+ }
536+ }
537+ 
538+ quantOutQueue.EnQue<Y>(quantOutLocal);
539+ xQueue.FreeTensor(xLocal);
540+}
541+} // namespace AdaLayerNormNS
542+ 
543+#endif // ADA_LAYER_NORM_WELFORD_H
Mnorm/ada_layer_norm/tests/ut/op_host/test_ada_layer_norm_tiling.cpp+8-156
@@ -1,10 +1,10 @@
1-/*1+/**
2- * This program is free software, you can redistribute it and/or modify.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3- * Copyright (c) 2025 Huawei Technologies Co., Ltd.3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4- * This file is a part of the CANN Open Software.4+ * CANN Open Software License Agreement Version 2.0 (the "License").
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.5 * 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 BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.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.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10#include <iostream>10#include <iostream>
@@ -57,6 +57,8 @@ TEST_F(AdaLayerNormTiling, AdaLayerNormTiling_001)
57 // compile info57 // compile info
58 struct AdaLayerNormCompileInfo {58 struct AdaLayerNormCompileInfo {
59 int32_t coreNum = 0;59 int32_t coreNum = 0;
60+ uint64_t ubSizePlatForm = 0;
61+ bool isRegBase = false;
60 } compile_info;62 } compile_info;
61 63 
62 std::string op_type("AdaLayerNorm");64 std::string op_type("AdaLayerNorm");
@@ -104,154 +106,4 @@ TEST_F(AdaLayerNormTiling, AdaLayerNormTiling_001)
104 // check tiling result106 // check tiling result
105 auto tiling_key = tiling_context->GetTilingKey();107 auto tiling_key = tiling_context->GetTilingKey();
106 ASSERT_EQ(tiling_key, 1);108 ASSERT_EQ(tiling_key, 1);
107-}
108- 
109-TEST_F(AdaLayerNormTiling, AdaLayerNormTiling_002)
110-{
111- string compile_info_string = R"({
112- "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
113- "Intrinsic_fix_pipe_l0c2out": false,
114- "Intrinsic_data_move_l12ub": true,
115- "Intrinsic_data_move_l0c2ub": true,
116- "Intrinsic_data_move_out2l1_nd2nz": false,
117- "UB_SIZE": 245760, "L2_SIZE": 33554432, "L1_SIZE": 524288,
118- "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
119- "CORE_NUM": 64}})";
120- map<string, string> soc_infos;
121- map<string, string> aicore_spec;
122- map<string, string> intrinsics;
123- GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
124- 
125- // platform info
126- fe::PlatFormInfos platform_info;
127- platform_info.Init();
128- // compile info
129- struct AdaLayerNormCompileInfo {
130- int32_t coreNum = 0;
131- } compile_info;
132- 
133- std::string op_type("AdaLayerNormQuant");
134- ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str()), nullptr);
135- auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
136- auto tiling_parse_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling_parse;
137- 
138- // tilingFunc simulate
139- auto param = gert::TilingData::CreateCap(4096);
140- ASSERT_NE(param, nullptr);
141- auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
142- auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
143- gert::StorageShape x = {{4, 16, 128}, {4, 16, 128}};
144- gert::StorageShape scale = {{4, 128}, {4, 128}};
145- gert::StorageShape shift = {{4, 128}, {4, 128}};
146- gert::StorageShape weight = {{128}, {128}};
147- gert::StorageShape bias = {{128}, {128}};
148- gert::StorageShape smoothScales = {{128}, {128}};
149- gert::StorageShape out = {{4, 16, 128}, {4, 16, 128}};
150- gert::StorageShape quantScale = {{4, 16}, {4, 16}};
151- auto holder = gert::TilingContextFaker()
152- .NodeIoNum(6, 2)
153- .IrInstanceNum({1, 1, 1, 1, 1, 1})
154- .InputShapes({&x, &scale, &shift, &weight, &bias, &smoothScales})
155- .OutputShapes({&out, &quantScale})
156- .CompileInfo(&compile_info)
157- .PlatformInfo(reinterpret_cast<char*>(&platform_info))
158- .NodeInputTd(0, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
159- .NodeInputTd(1, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
160- .NodeInputTd(2, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
161- .NodeInputTd(3, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
162- .NodeInputTd(4, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
163- .NodeInputTd(5, ge::DT_FLOAT16, ge::FORMAT_ND, ge::FORMAT_ND)
164- .NodeOutputTd(0, ge::DT_INT8, ge::FORMAT_ND, ge::FORMAT_ND)
165- .NodeOutputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
166- .NodeAttrs({{"epsilon", Ops::NN::AnyValue::CreateFrom<float>(0.00001)}})
167- .TilingData(param.get())
168- .Workspace(ws_size)
169- .Build();
170- 
171- gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
172- ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
173- holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
174- holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
175- holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
176- holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
177- 
178- EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
179- // check tiling result
180- auto tiling_key = tiling_context->GetTilingKey();
181- ASSERT_EQ(tiling_key, 1);
182-}
183- 
184-TEST_F(AdaLayerNormTiling, AdaLayerNormTiling_003)
185-{
186- string compile_info_string = R"({
187- "hardware_info": {"BT_SIZE": 0, "load3d_constraints": "1",
188- "Intrinsic_fix_pipe_l0c2out": false,
189- "Intrinsic_data_move_l12ub": true,
190- "Intrinsic_data_move_l0c2ub": true,
191- "Intrinsic_data_move_out2l1_nd2nz": false,
192- "UB_SIZE": 245760, "L2_SIZE": 33554432, "L1_SIZE": 524288,
193- "L0A_SIZE": 65536, "L0B_SIZE": 65536, "L0C_SIZE": 131072,
194- "CORE_NUM": 64}})";
195- map<string, string> soc_infos;
196- map<string, string> aicore_spec;
197- map<string, string> intrinsics;
198- GetPlatFormInfos(compile_info_string.c_str(), soc_infos, aicore_spec, intrinsics);
199- 
200- // platform info
201- fe::PlatFormInfos platform_info;
202- platform_info.Init();
203- // compile info
204- struct AdaLayerNormCompileInfo {
205- int32_t coreNum = 0;
206- } compile_info;
207- 
208- std::string op_type("AdaLayerNormV2");
209- ASSERT_NE(gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str()), nullptr);
210- auto tiling_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling;
211- auto tiling_parse_func = gert::OpImplRegistry::GetInstance().GetOpImpl(op_type.c_str())->tiling_parse;
212- 
213- // tilingFunc simulate
214- auto param = gert::TilingData::CreateCap(4096);
215- ASSERT_NE(param, nullptr);
216- auto workspace_size_holer = gert::ContinuousVector::Create<size_t>(4096);
217- auto ws_size = reinterpret_cast<gert::ContinuousVector*>(workspace_size_holer.get());
218- gert::StorageShape x = {{4, 16, 128}, {4, 16, 128}};
219- gert::StorageShape scale = {{4, 128}, {4, 128}};
220- gert::StorageShape shift = {{4, 128}, {4, 128}};
221- gert::StorageShape weight = {{128}, {128}};
222- gert::StorageShape bias = {{128}, {128}};
223- gert::StorageShape out = {{4, 16, 128}, {4, 16, 128}};
224- gert::StorageShape mean = {{4, 16, 1}, {4, 16, 1}};
225- gert::StorageShape rstd = {{4, 16, 1}, {4, 16, 1}};
226- auto holder = gert::TilingContextFaker()
227- .NodeIoNum(5, 3)
228- .IrInstanceNum({1, 1, 1, 1, 1})
229- .InputShapes({&x, &scale, &shift, &weight, &bias})
230- .OutputShapes({&out, &mean, &rstd})
231- .CompileInfo(&compile_info)
232- .PlatformInfo(reinterpret_cast<char*>(&platform_info))
233- .NodeInputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
234- .NodeInputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
235- .NodeInputTd(2, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
236- .NodeInputTd(3, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
237- .NodeInputTd(4, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
238- .NodeOutputTd(0, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
239- .NodeOutputTd(1, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
240- .NodeOutputTd(2, ge::DT_FLOAT, ge::FORMAT_ND, ge::FORMAT_ND)
241- .NodeAttrs({{"epsilon", Ops::NN::AnyValue::CreateFrom<float>(0.00001)}})
242- .TilingData(param.get())
243- .Workspace(ws_size)
244- .Build();
245- 
246- gert::TilingContext* tiling_context = holder.GetContext<gert::TilingContext>();
247- ASSERT_NE(tiling_context->GetPlatformInfo(), nullptr);
248- holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("SoCInfo", soc_infos);
249- holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreSpec", aicore_spec);
250- holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetCoreNumByCoreType("AICore");
251- holder.GetContext<gert::TilingContext>()->GetPlatformInfo()->SetPlatformRes("AICoreintrinsicDtypeMap", intrinsics);
252- 
253- EXPECT_EQ(tiling_func(tiling_context), ge::GRAPH_SUCCESS);
254- // check tiling result
255- auto tiling_key = tiling_context->GetTilingKey();
256- ASSERT_EQ(tiling_key, 2);
257}109}
Mnorm/ada_layer_norm_quant/op_host/ada_layer_norm_quant_def.cpp+51-42
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -9,59 +9,68 @@
9 */9 */
10 10 
11/*!11/*!
12- * \file ada_layer_norm_quant.cpp12+ * \file ada_layer_norm_quant_def.cpp
13 * \brief13 * \brief
14 */14 */
15#include "register/op_def_registry.h"15#include "register/op_def_registry.h"
16 16 
17namespace ops {17namespace ops {
18+static const std::vector<ge::DataType> xDataType = {ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_BF16,
19+ ge::DT_FLOAT16, ge::DT_BF16, ge::DT_FLOAT16, ge::DT_BF16};
20+static const std::vector<ge::DataType> outDataType = {ge::DT_INT8, ge::DT_INT8, ge::DT_HIFLOAT8,
21+ ge::DT_HIFLOAT8, ge::DT_FLOAT8_E4M3FN, ge::DT_FLOAT8_E4M3FN,
22+ ge::DT_FLOAT8_E5M2, ge::DT_FLOAT8_E5M2};
23+static const std::vector<ge::DataType> scaleDataType = {ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT,
24+ ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT, ge::DT_FLOAT};
25+static const std::vector<ge::Format> format = {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
CANN-robot
CANN-robotCANN-robot1月22日

代码结构与可维护性: format数组包含8个相同的FORMAT_ND格式。这种重复的硬编码同样属于'魔数'问题，且与数组长度8紧密耦合。如果未来配置数量变化，需要同步修改多个地方。

问题类型: 代码结构与可维护性 文件路径: norm/ada_layer_norm_quant/op_host/ada_layer_norm_quant_def.cpp 行号: 25 问题代码:

static const std::vector<ge::Format> format = {ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND,
                                               ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};

修改建议:

使用常量定义数组长度，并通过循环或std::fill初始化format数组。或者考虑是否真的需要8个相同的格式配置，或许可以简化为单个格式的配置。

此评论由代码审查工具自动生成

likedislike
26+ ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND};
27+ 
18class AdaLayerNormQuant : public OpDef {28class AdaLayerNormQuant : public OpDef {
19public:29public:
20 explicit AdaLayerNormQuant(const char *name) : OpDef(name)30 explicit AdaLayerNormQuant(const char *name) : OpDef(name)
21 {31 {
22- this->Input("x")32+ this->Input("x").ParamType(REQUIRED).DataType({ge::DT_FLOAT16, ge::DT_BF16})
23- .ParamType(REQUIRED)33+ .Format({ge::FORMAT_ND, ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
24- .DataType({ge::DT_FLOAT16, ge::DT_BF16})34+ this->Input("scale").ParamType(REQUIRED).DataType({ge::DT_FLOAT16, ge::DT_BF16})
25- .Format({ge::FORMAT_ND, ge::FORMAT_ND})35+ .Format({ge::FORMAT_ND, ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
26- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});36+ this->Input("shift").ParamType(REQUIRED).DataType({ge::DT_FLOAT16, ge::DT_BF16})
27- this->Input("scale")37+ .Format({ge::FORMAT_ND, ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
28- .ParamType(REQUIRED)38+ this->Input("weight").ParamType(OPTIONAL).DataType({ge::DT_FLOAT16, ge::DT_BF16})
29- .DataType({ge::DT_FLOAT16, ge::DT_BF16})39+ .Format({ge::FORMAT_ND, ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
30- .Format({ge::FORMAT_ND, ge::FORMAT_ND})40+ this->Input("bias").ParamType(OPTIONAL).DataType({ge::DT_FLOAT16, ge::DT_BF16})
31- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});41+ .Format({ge::FORMAT_ND, ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
32- this->Input("shift")42+ this->Input("smooth_scales").ParamType(OPTIONAL).DataType({ge::DT_FLOAT16, ge::DT_BF16})
33- .ParamType(REQUIRED)43+ .Format({ge::FORMAT_ND, ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
34- .DataType({ge::DT_FLOAT16, ge::DT_BF16})44+ this->Output("out").ParamType(REQUIRED).DataType({ge::DT_INT8, ge::DT_INT8})
35- .Format({ge::FORMAT_ND, ge::FORMAT_ND})45+ .Format({ge::FORMAT_ND, ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
36- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});46+ this->Output("quant_scale").ParamType(REQUIRED).DataType({ge::DT_FLOAT, ge::DT_FLOAT})
37- this->Input("weight")47+ .Format({ge::FORMAT_ND, ge::FORMAT_ND}).UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
38- .ParamType(OPTIONAL)
39- .DataType({ge::DT_FLOAT16, ge::DT_BF16})
40- .Format({ge::FORMAT_ND, ge::FORMAT_ND})
41- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
42- this->Input("bias")
43- .ParamType(OPTIONAL)
44- .DataType({ge::DT_FLOAT16, ge::DT_BF16})
45- .Format({ge::FORMAT_ND, ge::FORMAT_ND})
46- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
47- this->Input("smooth_scales")
48- .ParamType(OPTIONAL)
49- .DataType({ge::DT_FLOAT16, ge::DT_BF16})
50- .Format({ge::FORMAT_ND, ge::FORMAT_ND})
51- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
52- this->Output("out")
53- .ParamType(REQUIRED)
54- .DataType({ge::DT_INT8, ge::DT_INT8})
55- .Format({ge::FORMAT_ND, ge::FORMAT_ND})
56- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
57- this->Output("quant_scale")
58- .ParamType(REQUIRED)
59- .DataType({ge::DT_FLOAT, ge::DT_FLOAT})
60- .Format({ge::FORMAT_ND, ge::FORMAT_ND})
61- .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND});
62 this->Attr("epsilon").AttrType(OPTIONAL).Float(1e-5);48 this->Attr("epsilon").AttrType(OPTIONAL).Float(1e-5);
63 this->AICore().AddConfig("ascend910b");49 this->AICore().AddConfig("ascend910b");
64 this->AICore().AddConfig("ascend910_93");50 this->AICore().AddConfig("ascend910_93");
51+
52+ OpAICoreConfig config_91095;
53+ config_91095.Input("x").ParamType(REQUIRED).DataType(xDataType)
54+ .Format(format).UnknownShapeFormat(format);
55+ config_91095.Input("scale").ParamType(REQUIRED).DataType(xDataType)
56+ .Format(format).UnknownShapeFormat(format);
57+ config_91095.Input("shift").ParamType(REQUIRED).DataType(xDataType)
58+ .Format(format).UnknownShapeFormat(format);
59+ config_91095.Input("weight").ParamType(OPTIONAL).DataType(xDataType)
60+ .Format(format).UnknownShapeFormat(format);
61+ config_91095.Input("bias").ParamType(OPTIONAL).DataType(xDataType)
62+ .Format(format).UnknownShapeFormat(format);
63+ config_91095.Input("smooth_scales").ParamType(OPTIONAL).DataType(xDataType)
64+ .Format(format).UnknownShapeFormat(format);
65+ config_91095.Output("out").ParamType(REQUIRED).DataType(outDataType)
66+ .Format(format).UnknownShapeFormat(format);
67+ config_91095.Output("quant_scale").ParamType(REQUIRED).DataType(scaleDataType)
68+ .Format(format).UnknownShapeFormat(format);
69+ config_91095.DynamicCompileStaticFlag(true)
70+ .DynamicRankSupportFlag(true)
71+ .DynamicShapeSupportFlag(true)
72+ .ExtendCfgInfo("opFile.value", "ada_layer_norm_quant_apt");
73+ this->AICore().AddConfig("ascend910_95", config_91095);
65 }74 }
66};75};
67OP_ADD(AdaLayerNormQuant);76OP_ADD(AdaLayerNormQuant);
Anorm/ada_layer_norm_quant/op_host/config/ascend910_95/ada_layer_norm_quant_binary.json+757-0
@@ -0,0 +1,757 @@
1+{
2+ "op_type": "AdaLayerNormQuant",
3+ "op_list": [
4+ {
5+ "bin_filename": "AdaLayerNormQuant_8dddc26005b8c145f163a9c843f0de43",
6+ "inputs": [
7+ {
8+ "name": "x",
9+ "index": 0,
10+ "dtype": "float16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [
14+ -2
15+ ]
16+ },
17+ {
18+ "name": "scale",
19+ "index": 1,
20+ "dtype": "float16",
21+ "format": "ND",
22+ "paramType": "required",
23+ "shape": [
24+ -2
25+ ]
26+ },
27+ {
28+ "name": "shift",
29+ "index": 2,
30+ "dtype": "float16",
31+ "format": "ND",
32+ "paramType": "required",
33+ "shape": [
34+ -2
35+ ]
36+ },
37+ {
38+ "name": "weight",
39+ "index": 3,
40+ "dtype": "float16",
41+ "format": "ND",
42+ "paramType": "optional",
43+ "shape": [
44+ -2
45+ ]
46+ },
47+ {
48+ "name": "bias",
49+ "index": 4,
50+ "dtype": "float16",
51+ "format": "ND",
52+ "paramType": "optional",
53+ "shape": [
54+ -2
55+ ]
56+ },
57+ {
58+ "name": "smooth_scales",
59+ "index": 5,
60+ "dtype": "float16",
61+ "format": "ND",
62+ "paramType": "optional",
63+ "shape": [
64+ -2
65+ ]
66+ }
67+ ],
68+ "outputs": [
69+ {
70+ "name": "out",
71+ "index": 0,
72+ "dtype": "int8",
73+ "format": "ND",
74+ "paramType": "required",
75+ "shape": [
76+ -2
77+ ]
78+ },
79+ {
80+ "name": "quant_scale",
81+ "index": 1,
82+ "dtype": "float32",
83+ "format": "ND",
84+ "paramType": "required",
85+ "shape": [
86+ -2
87+ ]
88+ }
89+ ],
90+ "attrs": [
91+ {
92+ "name": "epsilon",
93+ "dtype": "float",
94+ "value": 0.0
95+ }
96+ ]
97+ },
98+ {
99+ "bin_filename": "AdaLayerNormQuant_38738e2582d51f99b5986cfe0b8b32f5",
100+ "inputs": [
101+ {
102+ "name": "x",
103+ "index": 0,
104+ "dtype": "bfloat16",
105+ "format": "ND",
106+ "paramType": "required",
107+ "shape": [
108+ -2
109+ ]
110+ },
111+ {
112+ "name": "scale",
113+ "index": 1,
114+ "dtype": "bfloat16",
115+ "format": "ND",
116+ "paramType": "required",
117+ "shape": [
118+ -2
119+ ]
120+ },
121+ {
122+ "name": "shift",
123+ "index": 2,
124+ "dtype": "bfloat16",
125+ "format": "ND",
126+ "paramType": "required",
127+ "shape": [
128+ -2
129+ ]
130+ },
131+ {
132+ "name": "weight",
133+ "index": 3,
134+ "dtype": "bfloat16",
135+ "format": "ND",
136+ "paramType": "optional",
137+ "shape": [
138+ -2
139+ ]
140+ },
141+ {
142+ "name": "bias",
143+ "index": 4,
144+ "dtype": "bfloat16",
145+ "format": "ND",
146+ "paramType": "optional",
147+ "shape": [
148+ -2
149+ ]
150+ },
151+ {
152+ "name": "smooth_scales",
153+ "index": 5,
154+ "dtype": "bfloat16",
155+ "format": "ND",
156+ "paramType": "optional",
157+ "shape": [
158+ -2
159+ ]
160+ }
161+ ],
162+ "outputs": [
163+ {
164+ "name": "out",
165+ "index": 0,
166+ "dtype": "int8",
167+ "format": "ND",
168+ "paramType": "required",
169+ "shape": [
170+ -2
171+ ]
172+ },
173+ {
174+ "name": "quant_scale",
175+ "index": 1,
176+ "dtype": "float32",
177+ "format": "ND",
178+ "paramType": "required",
179+ "shape": [
180+ -2
181+ ]
182+ }
183+ ],
184+ "attrs": [
185+ {
186+ "name": "epsilon",
187+ "dtype": "float",
188+ "value": 0.0
189+ }
190+ ]
191+ },
192+ {
193+ "bin_filename": "AdaLayerNormQuant_43f95645718722ece690db2a234ead4c",
194+ "inputs": [
195+ {
196+ "name": "x",
197+ "index": 0,
198+ "dtype": "float16",
199+ "format": "ND",
200+ "paramType": "required",
201+ "shape": [
202+ -2
203+ ]
204+ },
205+ {
206+ "name": "scale",
207+ "index": 1,
208+ "dtype": "float16",
209+ "format": "ND",
210+ "paramType": "required",
211+ "shape": [
212+ -2
213+ ]
214+ },
215+ {
216+ "name": "shift",
217+ "index": 2,
218+ "dtype": "float16",
219+ "format": "ND",
220+ "paramType": "required",
221+ "shape": [
222+ -2
223+ ]
224+ },
225+ {
226+ "name": "weight",
227+ "index": 3,
228+ "dtype": "float16",
229+ "format": "ND",
230+ "paramType": "optional",
231+ "shape": [
232+ -2
233+ ]
234+ },
235+ {
236+ "name": "bias",
237+ "index": 4,
238+ "dtype": "float16",
239+ "format": "ND",
240+ "paramType": "optional",
241+ "shape": [
242+ -2
243+ ]
244+ },
245+ {
246+ "name": "smooth_scales",
247+ "index": 5,
248+ "dtype": "float16",
249+ "format": "ND",
250+ "paramType": "optional",
251+ "shape": [
252+ -2
253+ ]
254+ }
255+ ],
256+ "outputs": [
257+ {
258+ "name": "out",
259+ "index": 0,
260+ "dtype": "float8_e5m2",
261+ "format": "ND",
262+ "paramType": "required",
263+ "shape": [
264+ -2
265+ ]
266+ },
267+ {
268+ "name": "quant_scale",
269+ "index": 1,
270+ "dtype": "float32",
271+ "format": "ND",
272+ "paramType": "required",
273+ "shape": [
274+ -2
275+ ]
276+ }
277+ ],
278+ "attrs": [
279+ {
280+ "name": "epsilon",
281+ "dtype": "float",
282+ "value": 0.0
283+ }
284+ ]
285+ },
286+ {
287+ "bin_filename": "AdaLayerNormQuant_f9e7c4de56a5a8145aac4fe95c1b373b",
288+ "inputs": [
289+ {
290+ "name": "x",
291+ "index": 0,
292+ "dtype": "bfloat16",
293+ "format": "ND",
294+ "paramType": "required",
295+ "shape": [
296+ -2
297+ ]
298+ },
299+ {
300+ "name": "scale",
301+ "index": 1,
302+ "dtype": "bfloat16",
303+ "format": "ND",
304+ "paramType": "required",
305+ "shape": [
306+ -2
307+ ]
308+ },
309+ {
310+ "name": "shift",
311+ "index": 2,
312+ "dtype": "bfloat16",
313+ "format": "ND",
314+ "paramType": "required",
315+ "shape": [
316+ -2
317+ ]
318+ },
319+ {
320+ "name": "weight",
321+ "index": 3,
322+ "dtype": "bfloat16",
323+ "format": "ND",
324+ "paramType": "optional",
325+ "shape": [
326+ -2
327+ ]
328+ },
329+ {
330+ "name": "bias",
331+ "index": 4,
332+ "dtype": "bfloat16",
333+ "format": "ND",
334+ "paramType": "optional",
335+ "shape": [
336+ -2
337+ ]
338+ },
339+ {
340+ "name": "smooth_scales",
341+ "index": 5,
342+ "dtype": "bfloat16",
343+ "format": "ND",
344+ "paramType": "optional",
345+ "shape": [
346+ -2
347+ ]
348+ }
349+ ],
350+ "outputs": [
351+ {
352+ "name": "out",
353+ "index": 0,
354+ "dtype": "float8_e5m2",
355+ "format": "ND",
356+ "paramType": "required",
357+ "shape": [
358+ -2
359+ ]
360+ },
361+ {
362+ "name": "quant_scale",
363+ "index": 1,
364+ "dtype": "float32",
365+ "format": "ND",
366+ "paramType": "required",
367+ "shape": [
368+ -2
369+ ]
370+ }
371+ ],
372+ "attrs": [
373+ {
374+ "name": "epsilon",
375+ "dtype": "float",
376+ "value": 0.0
377+ }
378+ ]
379+ },
380+ {
381+ "bin_filename": "AdaLayerNormQuant_e0d06064ace068ef0756f81b997cc2c2",
382+ "inputs": [
383+ {
384+ "name": "x",
385+ "index": 0,
386+ "dtype": "float16",
387+ "format": "ND",
388+ "paramType": "required",
389+ "shape": [
390+ -2
391+ ]
392+ },
393+ {
394+ "name": "scale",
395+ "index": 1,
396+ "dtype": "float16",
397+ "format": "ND",
398+ "paramType": "required",
399+ "shape": [
400+ -2
401+ ]
402+ },
403+ {
404+ "name": "shift",
405+ "index": 2,
406+ "dtype": "float16",
407+ "format": "ND",
408+ "paramType": "required",
409+ "shape": [
410+ -2
411+ ]
412+ },
413+ {
414+ "name": "weight",
415+ "index": 3,
416+ "dtype": "float16",
417+ "format": "ND",
418+ "paramType": "optional",
419+ "shape": [
420+ -2
421+ ]
422+ },
423+ {
424+ "name": "bias",
425+ "index": 4,
426+ "dtype": "float16",
427+ "format": "ND",
428+ "paramType": "optional",
429+ "shape": [
430+ -2
431+ ]
432+ },
433+ {
434+ "name": "smooth_scales",
435+ "index": 5,
436+ "dtype": "float16",
437+ "format": "ND",
438+ "paramType": "optional",
439+ "shape": [
440+ -2
441+ ]
442+ }
443+ ],
444+ "outputs": [
445+ {
446+ "name": "out",
447+ "index": 0,
448+ "dtype": "float8_e4m3fn",
449+ "format": "ND",
450+ "paramType": "required",
451+ "shape": [
452+ -2
453+ ]
454+ },
455+ {
456+ "name": "quant_scale",
457+ "index": 1,
458+ "dtype": "float32",
459+ "format": "ND",
460+ "paramType": "required",
461+ "shape": [
462+ -2
463+ ]
464+ }
465+ ],
466+ "attrs": [
467+ {
468+ "name": "epsilon",
469+ "dtype": "float",
470+ "value": 0.0
471+ }
472+ ]
473+ },
474+ {
475+ "bin_filename": "AdaLayerNormQuant_d96b1e780c7d122c24e56f348751ff72",
476+ "inputs": [
477+ {
478+ "name": "x",
479+ "index": 0,
480+ "dtype": "bfloat16",
481+ "format": "ND",
482+ "paramType": "required",
483+ "shape": [
484+ -2
485+ ]
486+ },
487+ {
488+ "name": "scale",
489+ "index": 1,
490+ "dtype": "bfloat16",
491+ "format": "ND",
492+ "paramType": "required",
493+ "shape": [
494+ -2
495+ ]
496+ },
497+ {
498+ "name": "shift",
499+ "index": 2,
500+ "dtype": "bfloat16",
501+ "format": "ND",
502+ "paramType": "required",
503+ "shape": [
504+ -2
505+ ]
506+ },
507+ {
508+ "name": "weight",
509+ "index": 3,
510+ "dtype": "bfloat16",
511+ "format": "ND",
512+ "paramType": "optional",
513+ "shape": [
514+ -2
515+ ]
516+ },
517+ {
518+ "name": "bias",
519+ "index": 4,
520+ "dtype": "bfloat16",
521+ "format": "ND",
522+ "paramType": "optional",
523+ "shape": [
524+ -2
525+ ]
526+ },
527+ {
528+ "name": "smooth_scales",
529+ "index": 5,
530+ "dtype": "bfloat16",
531+ "format": "ND",
532+ "paramType": "optional",
533+ "shape": [
534+ -2
535+ ]
536+ }
537+ ],
538+ "outputs": [
539+ {
540+ "name": "out",
541+ "index": 0,
542+ "dtype": "float8_e4m3fn",
543+ "format": "ND",
544+ "paramType": "required",
545+ "shape": [
546+ -2
547+ ]
548+ },
549+ {
550+ "name": "quant_scale",
551+ "index": 1,
552+ "dtype": "float32",
553+ "format": "ND",
554+ "paramType": "required",
555+ "shape": [
556+ -2
557+ ]
558+ }
559+ ],
560+ "attrs": [
561+ {
562+ "name": "epsilon",
563+ "dtype": "float",
564+ "value": 0.0
565+ }
566+ ]
567+ },
568+ {
569+ "bin_filename": "AdaLayerNormQuant_2504f2de6ad317fe8bf8eb443aba7cd0",
570+ "inputs": [
571+ {
572+ "name": "x",
573+ "index": 0,
574+ "dtype": "float16",
575+ "format": "ND",
576+ "paramType": "required",
577+ "shape": [
578+ -2
579+ ]
580+ },
581+ {
582+ "name": "scale",
583+ "index": 1,
584+ "dtype": "float16",
585+ "format": "ND",
586+ "paramType": "required",
587+ "shape": [
588+ -2
589+ ]
590+ },
591+ {
592+ "name": "shift",
593+ "index": 2,
594+ "dtype": "float16",
595+ "format": "ND",
596+ "paramType": "required",
597+ "shape": [
598+ -2
599+ ]
600+ },
601+ {
602+ "name": "weight",
603+ "index": 3,
604+ "dtype": "float16",
605+ "format": "ND",
606+ "paramType": "optional",
607+ "shape": [
608+ -2
609+ ]
610+ },
611+ {
612+ "name": "bias",
613+ "index": 4,
614+ "dtype": "float16",
615+ "format": "ND",
616+ "paramType": "optional",
617+ "shape": [
618+ -2
619+ ]
620+ },
621+ {
622+ "name": "smooth_scales",
623+ "index": 5,
624+ "dtype": "float16",
625+ "format": "ND",
626+ "paramType": "optional",
627+ "shape": [
628+ -2
629+ ]
630+ }
631+ ],
632+ "outputs": [
633+ {
634+ "name": "out",
635+ "index": 0,
636+ "dtype": "hifloat8",
637+ "format": "ND",
638+ "paramType": "required",
639+ "shape": [
640+ -2
641+ ]
642+ },
643+ {
644+ "name": "quant_scale",
645+ "index": 1,
646+ "dtype": "float32",
647+ "format": "ND",
648+ "paramType": "required",
649+ "shape": [
650+ -2
651+ ]
652+ }
653+ ],
654+ "attrs": [
655+ {
656+ "name": "epsilon",
657+ "dtype": "float",
658+ "value": 0.0
659+ }
660+ ]
661+ },
662+ {
663+ "bin_filename": "AdaLayerNormQuant_b503df464cdc5f8819bba1c7a4b65bfc",
664+ "inputs": [
665+ {
666+ "name": "x",
667+ "index": 0,
668+ "dtype": "bfloat16",
669+ "format": "ND",
670+ "paramType": "required",
671+ "shape": [
672+ -2
673+ ]
674+ },
675+ {
676+ "name": "scale",
677+ "index": 1,
678+ "dtype": "bfloat16",
679+ "format": "ND",
680+ "paramType": "required",
681+ "shape": [
682+ -2
683+ ]
684+ },
685+ {
686+ "name": "shift",
687+ "index": 2,
688+ "dtype": "bfloat16",
689+ "format": "ND",
690+ "paramType": "required",
691+ "shape": [
692+ -2
693+ ]
694+ },
695+ {
696+ "name": "weight",
697+ "index": 3,
698+ "dtype": "bfloat16",
699+ "format": "ND",
700+ "paramType": "optional",
701+ "shape": [
702+ -2
703+ ]
704+ },
705+ {
706+ "name": "bias",
707+ "index": 4,
708+ "dtype": "bfloat16",
709+ "format": "ND",
710+ "paramType": "optional",
711+ "shape": [
712+ -2
713+ ]
714+ },
715+ {
716+ "name": "smooth_scales",
717+ "index": 5,
718+ "dtype": "bfloat16",
719+ "format": "ND",
720+ "paramType": "optional",
721+ "shape": [
722+ -2
723+ ]
724+ }
725+ ],
726+ "outputs": [
727+ {
728+ "name": "out",
729+ "index": 0,
730+ "dtype": "hifloat8",
731+ "format": "ND",
732+ "paramType": "required",
733+ "shape": [
734+ -2
735+ ]
736+ },
737+ {
738+ "name": "quant_scale",
739+ "index": 1,
740+ "dtype": "float32",
741+ "format": "ND",
742+ "paramType": "required",
743+ "shape": [
744+ -2
745+ ]
746+ }
747+ ],
748+ "attrs": [
749+ {
750+ "name": "epsilon",
751+ "dtype": "float",
752+ "value": 0.0
753+ }
754+ ]
755+ }
756+ ]
757+}
Anorm/ada_layer_norm_quant/op_host/config/ascend910_95/ada_layer_norm_quant_simplified_key.ini+13-0
@@ -0,0 +1,13 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值，格式如下所示：
2+; [某算子]
3+; default=xx
4+; ascendxx=xx
5+; 其中，default为默认mode，ascendxx为可选mode，如果不同芯片有差异化要求时，需要配置；
6+; 1)如果没有配置：非ascendC算子继续按空处理，即opc编译命令中不添加 --simplified_key_mode 选项，AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置：各个版本按default配置
8+; 3)如果仅有某些平台的配置，没有default配置：对应平台的按照配置的值传递，非对应平台的：非AscendC算子继续按空处理，AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有：对应平台的使用平台的配置，非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况，需要在binary_simplified_key_mode.ini 文件中显式配置为None，不传入 --simplified_key_mode 选项，由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子，由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+[AdaLayerNormQuant]
13+default=0
Mnorm/ada_layer_norm_quant/op_host/op_api/aclnn_ada_layer_norm_quant.cpp+12-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -15,6 +15,7 @@
15#include "opdev/data_type_utils.h"15#include "opdev/data_type_utils.h"
16#include "opdev/format_utils.h"16#include "opdev/format_utils.h"
17#include "opdev/make_op_executor.h"17#include "opdev/make_op_executor.h"
18+#include "op_api/aclnn_util.h"
18#include "ada_layer_norm_quant.h"19#include "ada_layer_norm_quant.h"
19#include "aclnn_ada_layer_norm_quant.h"20#include "aclnn_ada_layer_norm_quant.h"
20 21 
@@ -47,6 +48,9 @@ static const std::initializer_list<op::DataType> X_DTYPE_SUPPORT_LIST = {
47 48 
48static const std::initializer_list<op::DataType> OUT_DTYPE_SUPPORT_LIST = {op::DataType::DT_INT8};49static const std::initializer_list<op::DataType> OUT_DTYPE_SUPPORT_LIST = {op::DataType::DT_INT8};
49 50 
51+static const std::initializer_list<op::DataType> OUT_DTYPE_SUPPORT_LIST_REGBASE = {
52+ op::DataType::DT_INT8, op::DataType::DT_HIFLOAT8, op::DataType::DT_FLOAT8_E5M2, op::DataType::DT_FLOAT8_E4M3FN};
53+ 
50static const std::initializer_list<op::DataType> SCALE_DTYPE_SUPPORT_LIST = {op::DataType::DT_FLOAT};54static const std::initializer_list<op::DataType> SCALE_DTYPE_SUPPORT_LIST = {op::DataType::DT_FLOAT};
51 55 
52static bool CheckNotNull(AdaLayerNormQuantInputTensor& inputTensor, AdaLayerNormQuantOutputTensor& outputTensor)56static bool CheckNotNull(AdaLayerNormQuantInputTensor& inputTensor, AdaLayerNormQuantOutputTensor& outputTensor)
@@ -73,7 +77,11 @@ static bool CheckDtypeValid(AdaLayerNormQuantInputTensor& inputTensor, AdaLayerN
73 if (inputTensor.smoothScalesOptional != nullptr) {77 if (inputTensor.smoothScalesOptional != nullptr) {
74 OP_CHECK_DTYPE_NOT_MATCH(inputTensor.smoothScalesOptional, inputTensor.x->GetDataType(), return false);78 OP_CHECK_DTYPE_NOT_MATCH(inputTensor.smoothScalesOptional, inputTensor.x->GetDataType(), return false);
75 }79 }
76- OP_CHECK_DTYPE_NOT_SUPPORT(outputTensor.out, OUT_DTYPE_SUPPORT_LIST, return false);80+ if (Ops::NN::AclnnUtil::IsRegbase()) {
81+ OP_CHECK_DTYPE_NOT_SUPPORT(outputTensor.out, OUT_DTYPE_SUPPORT_LIST_REGBASE, return false);
82+ } else {
83+ OP_CHECK_DTYPE_NOT_SUPPORT(outputTensor.out, OUT_DTYPE_SUPPORT_LIST, return false);
84+ }
77 OP_CHECK_DTYPE_NOT_SUPPORT(outputTensor.quantScale, SCALE_DTYPE_SUPPORT_LIST, return false);85 OP_CHECK_DTYPE_NOT_SUPPORT(outputTensor.quantScale, SCALE_DTYPE_SUPPORT_LIST, return false);
78 if (outputTensor.quantOffsetOptional != nullptr) {86 if (outputTensor.quantOffsetOptional != nullptr) {
79 OP_LOGE(ACLNN_ERR_PARAM_INVALID, "quantOffsetOptional should be nullptr.");87 OP_LOGE(ACLNN_ERR_PARAM_INVALID, "quantOffsetOptional should be nullptr.");
@@ -208,9 +216,10 @@ aclnnStatus aclnnAdaLayerNormQuantGetWorkspaceSize(
208 CHECK_RET(smoothScalesOptional != nullptr, ACLNN_ERR_INNER_NULLPTR);216 CHECK_RET(smoothScalesOptional != nullptr, ACLNN_ERR_INNER_NULLPTR);
209 }217 }
210 218 
219+ int32_t dstType = out->GetDataType();
211 std::tuple<aclTensor*, aclTensor*> result = l0op::AdaLayerNormQuant(220 std::tuple<aclTensor*, aclTensor*> result = l0op::AdaLayerNormQuant(
212 x, scale, shift, weightOptional, biasOptional, smoothScalesOptional, static_cast<float>(epsilon), quantMode,221 x, scale, shift, weightOptional, biasOptional, smoothScalesOptional, static_cast<float>(epsilon), quantMode,
213- uniqueExecutor.get());222+ dstType, uniqueExecutor.get());
214 const aclTensor* resultTensor = std::get<0>(result);223 const aclTensor* resultTensor = std::get<0>(result);
215 const aclTensor* quantScaleTensor = std::get<1>(result);224 const aclTensor* quantScaleTensor = std::get<1>(result);
216 CHECK_RET(resultTensor != nullptr && quantScaleTensor != nullptr, ACLNN_ERR_INNER_NULLPTR);225 CHECK_RET(resultTensor != nullptr && quantScaleTensor != nullptr, ACLNN_ERR_INNER_NULLPTR);
Mnorm/ada_layer_norm_quant/op_host/op_api/ada_layer_norm_quant.cpp+3-3
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -27,7 +27,7 @@ OP_TYPE_REGISTER(AdaLayerNormQuant);
27const std::tuple<aclTensor*, aclTensor*> AdaLayerNormQuant(27const std::tuple<aclTensor*, aclTensor*> AdaLayerNormQuant(
28 const aclTensor* x, const aclTensor* scale, const aclTensor* shift, const aclTensor* weightOptional,28 const aclTensor* x, const aclTensor* scale, const aclTensor* shift, const aclTensor* weightOptional,
29 const aclTensor* biasOptional, const aclTensor* smoothScalesOptional, float epsilon, const char* quantMode,29 const aclTensor* biasOptional, const aclTensor* smoothScalesOptional, float epsilon, const char* quantMode,
30- aclOpExecutor* executor)30+ int32_t dstType, aclOpExecutor* executor)
31{31{
32 L0_DFX(AdaLayerNormQuant, x, scale, shift, weightOptional, biasOptional, smoothScalesOptional, epsilon, quantMode);32 L0_DFX(AdaLayerNormQuant, x, scale, shift, weightOptional, biasOptional, smoothScalesOptional, epsilon, quantMode);
33 33 
@@ -36,7 +36,7 @@ const std::tuple<aclTensor*, aclTensor*> AdaLayerNormQuant(
36 for (size_t i = 0; i < dimNum - 1; i++) {36 for (size_t i = 0; i < dimNum - 1; i++) {
37 scaleShape.AppendDim(x->GetViewShape().GetDim(i));37 scaleShape.AppendDim(x->GetViewShape().GetDim(i));
38 }38 }
39- auto out = executor->AllocTensor(x->GetViewShape(), DataType::DT_INT8, x->GetViewFormat());39+ auto out = executor->AllocTensor(x->GetViewShape(), op::DataType(dstType), x->GetViewFormat());
40 auto quantScale = executor->AllocTensor(scaleShape, DataType::DT_FLOAT, op::Format::FORMAT_ND);40 auto quantScale = executor->AllocTensor(scaleShape, DataType::DT_FLOAT, op::Format::FORMAT_ND);
41 41 
42 auto ret = ADD_TO_LAUNCHER_LIST_AICORE(42 auto ret = ADD_TO_LAUNCHER_LIST_AICORE(
Mnorm/ada_layer_norm_quant/op_host/op_api/ada_layer_norm_quant.h+2-2
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -17,7 +17,7 @@ namespace l0op {
17const std::tuple<aclTensor*, aclTensor*> AdaLayerNormQuant(17const std::tuple<aclTensor*, aclTensor*> AdaLayerNormQuant(
18 const aclTensor* x, const aclTensor* scale, const aclTensor* shift, const aclTensor* weightOptional,18 const aclTensor* x, const aclTensor* scale, const aclTensor* shift, const aclTensor* weightOptional,
19 const aclTensor* biasOptional, const aclTensor* smoothScalesOptional, float epsilon, const char* quantMode,19 const aclTensor* biasOptional, const aclTensor* smoothScalesOptional, float epsilon, const char* quantMode,
20- aclOpExecutor* executor);20+ int32_t dstType, aclOpExecutor* executor);
21} // namespace l0op21} // namespace l0op
22 22 
23#endif // OP_API_INC_LEVEL0_ADA_LAYER_NORM_QUANT_H23#endif // OP_API_INC_LEVEL0_ADA_LAYER_NORM_QUANT_H
Mnorm/ada_layer_norm_quant/op_kernel/ada_layer_norm_quant.cpp+2-1
@@ -1,5 +1,5 @@
1/**1/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.2+ * Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.
3 * This program is free software, you can redistribute it and/or modify it under the terms and conditions of3 * 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").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.5 * Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -22,6 +22,7 @@ extern "C" __global__ __aicore__ void ada_layer_norm_quant(
22 GM_ADDR x, GM_ADDR scale, GM_ADDR shift, GM_ADDR weight, GM_ADDR bias, GM_ADDR smooth_scales, GM_ADDR out,22 GM_ADDR x, GM_ADDR scale, GM_ADDR shift, GM_ADDR weight, GM_ADDR bias, GM_ADDR smooth_scales, GM_ADDR out,
23 GM_ADDR quant_scale, GM_ADDR workspace, GM_ADDR tiling)23 GM_ADDR quant_scale, GM_ADDR workspace, GM_ADDR tiling)
24{24{
25+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
25 GET_TILING_DATA(tilingData, tiling);26 GET_TILING_DATA(tilingData, tiling);
26 27 
27 GM_ADDR userWS = GetUserWorkspace(workspace);28 GM_ADDR userWS = GetUserWorkspace(workspace);
Anorm/ada_layer_norm_quant/op_kernel/ada_layer_norm_quant_apt.cpp+44-0
@@ -0,0 +1,44 @@
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 ada_layer_norm_quant_apt.cpp
13+ * \brief
14+ */
15+ 
16+#include "../ada_layer_norm/arch35/ada_layer_norm_impl.h"
17+ 
18+using namespace AdaLayerNormNS;
19+ 
20+extern "C" __global__ __aicore__ void ada_layer_norm_quant(
21+ GM_ADDR x, GM_ADDR scale, GM_ADDR shift, GM_ADDR weight, GM_ADDR bias, GM_ADDR smooth_scales, GM_ADDR out,
22+ GM_ADDR quant_scale, GM_ADDR workspace, GM_ADDR tiling)
23+{
24+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
25+ GET_TILING_DATA(tilingData, tiling);
26+ 
27+ GM_ADDR usrWorkspace = GetUserWorkspace(workspace);
28+ if (usrWorkspace == nullptr) {
29+ return;
30+ }
31+ 
32+#define INIT_AND_PROCESS \
33+ op.InitQuant(&gmAddr, usrWorkspace, &tilingData); \
34+ op.Process()
35+ 
36+ GmAddr gmAddr = {x, scale, shift, weight, bias, smooth_scales, out, nullptr, nullptr, quant_scale};
37+ if (TILING_KEY_IS(11)) {
38+ AdaLayerNormFullLoad<DTYPE_X, DTYPE_X, DTYPE_OUT, QUANT_OP_CODE> op;
39+ INIT_AND_PROCESS;
40+ } else if (TILING_KEY_IS(21)) {
41+ AdaLayerNormWelford<DTYPE_X, DTYPE_X, DTYPE_OUT, QUANT_OP_CODE> op;
42+ INIT_AND_PROCESS;
43+ }
44+}
Mnorm/ada_layer_norm_quant/tests/ut/op_host/CMakeLists.txt+1-1
@@ -9,7 +9,7 @@
9 9 
10file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)10file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
11if(UT_TEST_ALL OR OP_HOST_UT)11if(UT_TEST_ALL OR OP_HOST_UT)
12- #add_modules_ut_sources(HOSTNAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})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})13 #add_modules_ut_sources(HOSTNAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14endif()14endif()
15if(UT_TEST_ALL OR OP_API_UT)15if(UT_TEST_ALL OR OP_API_UT)