Star1043
1939
代码
代码
Issues164
Pull Requests310
Actions
项目讨论
Wiki
分析
Star1043
1939
已合并
fix(foreach): AddListV2 整数溢出改为二进制补码回绕语义 #9151
Tian_1122创建于 8月25日
fix(foreach): AddListV2 整数溢出改为二进制补码回绕语义 #9151
已合并
Tian_1122创建于 8月25日
从Tian_1122/ops-nnmaster合入到cann/ops-nnmaster
共 2 个文件变更+287-3
Mforeach/foreach_add_list/op_kernel/foreach_add_list.cpp+4-3
@@ -18,6 +18,7 @@
18// op kernel building at build_out directory, it's not fully aligned with source code structure18// op kernel building at build_out directory, it's not fully aligned with source code structure
19// current op_kernel folder is absent in build_out directory, so the relative path to common has just one layer19// current op_kernel folder is absent in build_out directory, so the relative path to common has just one layer
20#include "../foreach_utils/foreach_one_scalar_ternary.h"20#include "../foreach_utils/foreach_one_scalar_ternary.h"
21+#include "foreach_add_list_wrap.h"
21 22 
22using namespace AscendC;23using namespace AscendC;
23using namespace Common::OpKernel;24using namespace Common::OpKernel;
@@ -77,15 +78,15 @@ extern "C" __global__ __aicore__ void foreach_add_list(GM_ADDR inputs_1, GM_ADDR
77 op.Init(inputs_1, inputs_2, alpha, outputs, userWS, &tilingData);78 op.Init(inputs_1, inputs_2, alpha, outputs, userWS, &tilingData);
78 op.Process();79 op.Process();
79 } else if (TILING_KEY_IS(5)) {80 } else if (TILING_KEY_IS(5)) {
80- ForeachOneScalarTernary<int16_t, float, AddListFloatAdapter<float>> op;81+ ForeachAddListWrap<int16_t> op;
81 op.Init(inputs_1, inputs_2, alpha, outputs, userWS, &tilingData);82 op.Init(inputs_1, inputs_2, alpha, outputs, userWS, &tilingData);
82 op.Process();83 op.Process();
83 } else if (TILING_KEY_IS(7)) {84 } else if (TILING_KEY_IS(7)) {
84- ForeachOneScalarTernary<int8_t, half, AddListFloatAdapter<half>> op;85+ ForeachAddListWrap<int8_t> op;
85 op.Init(inputs_1, inputs_2, alpha, outputs, userWS, &tilingData);86 op.Init(inputs_1, inputs_2, alpha, outputs, userWS, &tilingData);
86 op.Process();87 op.Process();
87 } else if (TILING_KEY_IS(8)) {88 } else if (TILING_KEY_IS(8)) {
88- ForeachOneScalarTernary<uint8_t, half, AddListFloatAdapter<half>> op;89+ ForeachAddListWrap<uint8_t> op;
89 op.Init(inputs_1, inputs_2, alpha, outputs, userWS, &tilingData);90 op.Init(inputs_1, inputs_2, alpha, outputs, userWS, &tilingData);
90 op.Process();91 op.Process();
91#endif92#endif
Aforeach/foreach_add_list/op_kernel/foreach_add_list_wrap.h+283-0
@@ -0,0 +1,283 @@
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 foreach_add_list_wrap.h
13+ * \brief foreach_add_list int16/int8/uint8 kernel, 整数溢出按二进制补码回绕
14+ */
15+ 
16+#ifndef FOREACH_ADD_LIST_WRAP_H
17+#define FOREACH_ADD_LIST_WRAP_H
18+ 
19+#include "kernel_operator.h"
20+ 
21+namespace Common {
22+namespace OpKernel {
23+using namespace AscendC;
24+ 
25+constexpr int32_t ADD_LIST_WRAP_BUFFER_NUM = 2;
26+ 
27+template <typename T, int32_t bufferNum = ADD_LIST_WRAP_BUFFER_NUM>
28+class ForeachAddListWrap {
29+public:
30+ __aicore__ inline ForeachAddListWrap(){};
31+ __aicore__ inline void Init(GM_ADDR x1, GM_ADDR x2, GM_ADDR alpha, GM_ADDR y, GM_ADDR workspace,
32+ const ForeachCommonTilingData* tilingData);
33+ __aicore__ inline void Process();
34+ 
35+private:
36+ __aicore__ inline void ParseTilingData(const ForeachCommonTilingData* tilingData);
37+ __aicore__ inline void SingleTensorProcess(int64_t dataCount);
38+ __aicore__ inline void CopyIn(uint32_t index, int64_t dataCount, bool isRemainder);
39+ __aicore__ inline void CopyIn2(uint32_t index, int64_t dataCount, bool isRemainder);
40+ __aicore__ inline void Compute(uint32_t index, int64_t dataCount, bool isRemainder);
41+ __aicore__ inline void CopyOut(uint32_t index, int64_t dataCount, bool isRemainder);
42+ __aicore__ inline __gm__ T* GetTensorAddr(uint16_t index, GM_ADDR tensorPtr);
43+ 
44+private:
45+ TPipe pipe;
46+ TQue<QuePosition::VECIN, bufferNum> dataQueue1;
47+ TQue<QuePosition::VECIN, bufferNum> dataQueue2;
48+ TQue<QuePosition::VECOUT, bufferNum> outQueue;
49+ TBuf<QuePosition::VECCALC> halfBuf;
50+ TBuf<QuePosition::VECCALC> int16Buf;
51+ 
52+ GlobalTensor<T> inTensorsGM1;
53+ GlobalTensor<T> inTensorsGM2;
54+ GlobalTensor<T> outTensorsGM;
55+ GlobalTensor<DTYPE_ALPHA> inScalarGM;
56+ 
57+ GM_ADDR inTensorsPtr1 = nullptr;
58+ GM_ADDR inTensorsPtr2 = nullptr;
59+ GM_ADDR outTensorsPtr = nullptr;
60+ 
61+ int16_t alphaVal = 0;
62+ 
63+ int64_t blockIdx = 0;
64+ uint32_t maxDataCount = 0;
65+ uint64_t inputsTensorUbSize = 0;
66+ const int64_t* tensorDataCountList = nullptr;
67+ uint16_t tensorStart = 0;
68+ uint16_t tensorEnd = 0;
69+ int64_t tensorStartOffset = 0;
70+ int64_t tensorEndOffset = 0;
71+};
72+ 
73+template <typename T, int32_t bufferNum>
74+__aicore__ inline void ForeachAddListWrap<T, bufferNum>::Init(GM_ADDR x1, GM_ADDR x2, GM_ADDR alpha, GM_ADDR y,
75+ GM_ADDR workspace,
76+ const ForeachCommonTilingData* tilingData)
77+{
78+ (void)workspace;
79+ blockIdx = GetBlockIdx();
80+ inTensorsPtr1 = x1;
81+ inTensorsPtr2 = x2;
82+ outTensorsPtr = y;
83+ ParseTilingData(tilingData);
84+ 
85+ inScalarGM.SetGlobalBuffer((__gm__ DTYPE_ALPHA*)alpha, 1);
86+ alphaVal = static_cast<int16_t>(inScalarGM.GetValue(0));
87+ 
88+ if constexpr (std::is_same_v<T, int8_t> || std::is_same_v<T, uint8_t>) {
89+ // int8/uint8: 经 half 提升 int16 计算, 再按位掩码回绕到目标 dtype 低比特
90+ maxDataCount = static_cast<uint32_t>(inputsTensorUbSize);
91+ pipe.InitBuffer(dataQueue1, bufferNum, maxDataCount * sizeof(T));
92+ pipe.InitBuffer(dataQueue2, bufferNum, maxDataCount * sizeof(T));
93+ pipe.InitBuffer(outQueue, bufferNum, maxDataCount * sizeof(T));
94+ pipe.InitBuffer(halfBuf, ADD_LIST_WRAP_BUFFER_NUM * maxDataCount * sizeof(half));
95+ // [a16 | b16 | mask | lo | hi]
96+ pipe.InitBuffer(int16Buf, 5 * maxDataCount * sizeof(int16_t));
97+ } else {
98+ // int16: 直接 int16 域计算, 硬件按二进制补码回绕, 无需中间转换
99+ maxDataCount = static_cast<uint32_t>(inputsTensorUbSize / sizeof(T));
100+ pipe.InitBuffer(dataQueue1, bufferNum, maxDataCount * sizeof(T));
101+ pipe.InitBuffer(dataQueue2, bufferNum, maxDataCount * sizeof(T));
102+ pipe.InitBuffer(outQueue, bufferNum, maxDataCount * sizeof(T));
103+ }
104+}
105+ 
106+template <typename T, int32_t bufferNum>
107+__aicore__ inline void ForeachAddListWrap<T, bufferNum>::ParseTilingData(const ForeachCommonTilingData* tilingData)
108+{
109+ inputsTensorUbSize = tilingData->inputsTensorUbSize;
110+ tensorDataCountList = tilingData->tensorDataCountList;
111+ tensorStart = tilingData->tensorStartList[blockIdx];
112+ tensorEnd = tilingData->tensorEndList[blockIdx];
113+ tensorStartOffset = tilingData->tensorStartOffsetList[blockIdx];
114+ tensorEndOffset = tilingData->tensorEndOffsetList[blockIdx];
115+}
116+ 
117+template <typename T, int32_t bufferNum>
118+__aicore__ inline void ForeachAddListWrap<T, bufferNum>::Process()
119+{
120+ for (uint16_t i = tensorStart; i <= tensorEnd; i++) {
121+ int64_t cursorStart = 0;
122+ int64_t cursorEnd = tensorDataCountList[i] - 1;
123+ if (i == tensorStart) {
124+ cursorStart = tensorStartOffset;
125+ }
126+ if (i == tensorEnd) {
127+ cursorEnd = tensorEndOffset;
128+ }
129+ 
130+ int64_t dataCount = cursorEnd - cursorStart + 1;
131+ inTensorsGM1.SetGlobalBuffer(GetTensorAddr(i, inTensorsPtr1) + cursorStart);
132+ inTensorsGM2.SetGlobalBuffer(GetTensorAddr(i, inTensorsPtr2) + cursorStart);
133+ outTensorsGM.SetGlobalBuffer(GetTensorAddr(i, outTensorsPtr) + cursorStart);
134+ SingleTensorProcess(dataCount);
135+ }
136+}
137+ 
138+template <typename T, int32_t bufferNum>
139+__aicore__ inline void ForeachAddListWrap<T, bufferNum>::SingleTensorProcess(int64_t dataCount)
140+{
141+ uint32_t copyTimes = static_cast<uint32_t>(dataCount / maxDataCount);
142+ uint32_t copyTimesRemainder = static_cast<uint32_t>(dataCount % maxDataCount);
143+ uint32_t tempDataCount = maxDataCount;
144+ 
145+ if (copyTimesRemainder > 0) {
146+ copyTimes++;
147+ }
148+ 
149+ for (uint32_t i = 0; i < copyTimes; i++) {
150+ bool isRemainder = false;
151+ if (i == copyTimes - 1 && copyTimesRemainder > 0) {
152+ isRemainder = true;
153+ tempDataCount = copyTimesRemainder;
154+ }
155+ CopyIn(i, tempDataCount, isRemainder);
156+ CopyIn2(i, tempDataCount, isRemainder);
157+ Compute(i, tempDataCount, isRemainder);
158+ CopyOut(i, tempDataCount, isRemainder);
159+ }
160+}
161+ 
162+template <typename T, int32_t bufferNum>
163+__aicore__ inline void ForeachAddListWrap<T, bufferNum>::CopyIn(uint32_t index, int64_t dataCount, bool isRemainder)
164+{
165+ LocalTensor<T> dataLocal = dataQueue1.template AllocTensor<T>();
166+ if (isRemainder) {
167+ DataCopyExtParams copyParams{1, static_cast<uint32_t>(dataCount * sizeof(T)), 0, 0, 0};
168+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
169+ DataCopyPad(dataLocal, inTensorsGM1[1ULL * index * maxDataCount], copyParams, padParams);
170+ } else {
171+ DataCopy(dataLocal, inTensorsGM1[1ULL * index * maxDataCount], dataCount);
172+ }
173+ dataQueue1.EnQue(dataLocal);
174+}
175+ 
176+template <typename T, int32_t bufferNum>
177+__aicore__ inline void ForeachAddListWrap<T, bufferNum>::CopyIn2(uint32_t index, int64_t dataCount, bool isRemainder)
178+{
179+ LocalTensor<T> dataLocal = dataQueue2.template AllocTensor<T>();
180+ if (isRemainder) {
181+ DataCopyExtParams copyParams{1, static_cast<uint32_t>(dataCount * sizeof(T)), 0, 0, 0};
182+ DataCopyPadExtParams<T> padParams{false, 0, 0, 0};
183+ DataCopyPad(dataLocal, inTensorsGM2[1ULL * index * maxDataCount], copyParams, padParams);
184+ } else {
185+ DataCopy(dataLocal, inTensorsGM2[1ULL * index * maxDataCount], dataCount);
186+ }
187+ dataQueue2.EnQue(dataLocal);
188+}
189+ 
190+template <typename T, int32_t bufferNum>
191+__aicore__ inline void ForeachAddListWrap<T, bufferNum>::Compute(uint32_t index, int64_t dataCount, bool isRemainder)
192+{
193+ (void)index;
194+ (void)isRemainder;
195+ LocalTensor<T> inLocal1 = dataQueue1.template DeQue<T>();
196+ LocalTensor<T> inLocal2 = dataQueue2.template DeQue<T>();
197+ LocalTensor<T> outLocal = outQueue.template AllocTensor<T>();
198+ 
199+ PipeBarrier<PIPE_V>();
200+ if constexpr (std::is_same_v<T, int16_t>) {
201+ // int16 域: a + alpha * b, 溢出按二进制补码回绕
202+ Muls(inLocal2, inLocal2, alphaVal, dataCount);
203+ PipeBarrier<PIPE_V>();
204+ Add(outLocal, inLocal1, inLocal2, dataCount);
205+ PipeBarrier<PIPE_V>();
206+ } else {
207+ // int8/uint8: 提升 int16 精确计算后取低 8 比特
208+ LocalTensor<half> h1 = halfBuf.GetWithOffset<half>(maxDataCount, 0);
209+ LocalTensor<half> h2 = halfBuf.GetWithOffset<half>(maxDataCount, maxDataCount * sizeof(half));
210+ LocalTensor<int16_t> a16 = int16Buf.GetWithOffset<int16_t>(maxDataCount, 0);
211+ LocalTensor<int16_t> b16 = int16Buf.GetWithOffset<int16_t>(maxDataCount, maxDataCount * sizeof(int16_t));
212+ LocalTensor<int16_t> mask16 = int16Buf.GetWithOffset<int16_t>(maxDataCount, 2 * maxDataCount * sizeof(int16_t));
213+ LocalTensor<int16_t> lo16 = int16Buf.GetWithOffset<int16_t>(maxDataCount, 3 * maxDataCount * sizeof(int16_t));
214+ LocalTensor<int16_t> hi16 = int16Buf.GetWithOffset<int16_t>(maxDataCount, 4 * maxDataCount * sizeof(int16_t));
215+ 
216+ Cast(h1, inLocal1, RoundMode::CAST_NONE, dataCount);
217+ PipeBarrier<PIPE_V>();
218+ Cast(h2, inLocal2, RoundMode::CAST_NONE, dataCount);
219+ PipeBarrier<PIPE_V>();
220+ Cast(a16, h1, RoundMode::CAST_RINT, dataCount);
221+ PipeBarrier<PIPE_V>();
222+ Cast(b16, h2, RoundMode::CAST_RINT, dataCount);
223+ PipeBarrier<PIPE_V>();
224+ Muls(b16, b16, alphaVal, dataCount);
225+ PipeBarrier<PIPE_V>();
226+ Add(a16, a16, b16, dataCount);
227+ PipeBarrier<PIPE_V>();
228+ if constexpr (std::is_same_v<T, uint8_t>) {
229+ // 无符号回绕: 保留低 8 位
230+ Duplicate(mask16, static_cast<int16_t>(0x00FF), dataCount);
231+ PipeBarrier<PIPE_V>();
232+ And(a16, a16, mask16, dataCount);
233+ PipeBarrier<PIPE_V>();
234+ } else {
235+ // 有符号回绕: (v & 0x7F) - (v & 0x80), 映射到 [-128, 127]
236+ Duplicate(mask16, static_cast<int16_t>(0x007F), dataCount);
237+ PipeBarrier<PIPE_V>();
238+ And(lo16, a16, mask16, dataCount);
239+ PipeBarrier<PIPE_V>();
240+ Duplicate(mask16, static_cast<int16_t>(0x0080), dataCount);
241+ PipeBarrier<PIPE_V>();
242+ And(hi16, a16, mask16, dataCount);
243+ PipeBarrier<PIPE_V>();
244+ Sub(a16, lo16, hi16, dataCount);
245+ PipeBarrier<PIPE_V>();
246+ }
247+ Cast(h1, a16, RoundMode::CAST_NONE, dataCount);
248+ PipeBarrier<PIPE_V>();
249+ Cast(outLocal, h1, RoundMode::CAST_RINT, dataCount);
250+ PipeBarrier<PIPE_V>();
251+ }
252+ 
253+ outQueue.EnQue(outLocal);
254+ dataQueue1.FreeTensor(inLocal1);
255+ dataQueue2.FreeTensor(inLocal2);
256+}
257+ 
258+template <typename T, int32_t bufferNum>
259+__aicore__ inline void ForeachAddListWrap<T, bufferNum>::CopyOut(uint32_t index, int64_t dataCount, bool isRemainder)
260+{
261+ LocalTensor<T> outLocal = outQueue.template DeQue<T>();
262+ if (isRemainder) {
263+ DataCopyExtParams copyParams{1, static_cast<uint32_t>(dataCount * sizeof(T)), 0, 0, 0};
264+ DataCopyPad(outTensorsGM[1ULL * index * maxDataCount], outLocal, copyParams);
265+ } else {
266+ DataCopy(outTensorsGM[1ULL * index * maxDataCount], outLocal, dataCount);
267+ }
268+ outQueue.FreeTensor(outLocal);
269+}
270+ 
271+template <typename T, int32_t bufferNum>
272+__aicore__ inline __gm__ T* ForeachAddListWrap<T, bufferNum>::GetTensorAddr(uint16_t index, GM_ADDR tensorPtr)
273+{
274+ __gm__ uint64_t* dataAddr = reinterpret_cast<__gm__ uint64_t*>(tensorPtr);
275+ uint64_t tensorPtrOffset = *dataAddr;
276+ __gm__ uint64_t* retPtr = dataAddr + (tensorPtrOffset >> 3);
277+ return reinterpret_cast<__gm__ T*>(*(retPtr + index));
278+}
279+ 
280+} // namespace OpKernel
281+} // namespace Common
282+ 
283+#endif // FOREACH_ADD_LIST_WRAP_H