Star872
1766
代码
代码
Issues242
Pull Requests323
Actions
项目讨论
Wiki
分析
Star872
1766
已合并
fix foreach_addcmul_scalar int32 经 float24 往返丢失低位 #8650
east_yang创建于 13 天前
fix foreach_addcmul_scalar int32 经 float24 往返丢失低位 #8650
已合并
east_yang创建于 13 天前
从east_yang/ops-nnmaster合入到cann/ops-nnmaster
共 5 个文件变更+165-141
Mactivation/selu_grad/examples/test_aclnn_selu_grad.cpp+127-122
@@ -13,141 +13,146 @@
13#include "aclnnop/aclnn_selu_backward.h"13#include "aclnnop/aclnn_selu_backward.h"
14 14 
15#define CHECK_RET(cond, return_expr) \15#define CHECK_RET(cond, return_expr) \
16- do { \16+ do { \
17- if (!(cond)) { \17+ if (!(cond)) { \
18- return_expr; \18+ return_expr; \
19- } \19+ } \
20- } while(0)20+ } while (0)
21 21 
22-#define LOG_PRINT(message, ...) \22+#define LOG_PRINT(message, ...) \
23- do { \23+ do { \
24- printf(message, ##__VA_ARGS__); \24+ printf(message, ##__VA_ARGS__); \
25- } while(0)25+ } while (0)
26 26 
27-int64_t GetShapeSize(const std::vector<int64_t>& shape) {27+int64_t GetShapeSize(const std::vector<int64_t>& shape)
28- int64_t shape_size = 1;28+{
29- for (auto i : shape) {29+ int64_t shape_size = 1;
30- shape_size *= i;30+ for (auto i : shape) {
31- }31+ shape_size *= i;
32- return shape_size;32+ }
33+ return shape_size;
33}34}
34 35 
35-int Init(int32_t deviceId, aclrtStream* stream) {36+int Init(int32_t deviceId, aclrtStream* stream)
36- // 固定写法,资源初始化37+{
37- auto ret = aclInit(nullptr);38+ // 固定写法,资源初始化
38- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);39+ auto ret = aclInit(nullptr);
39- ret = aclrtSetDevice(deviceId);40+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
40- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);41+ ret = aclrtSetDevice(deviceId);
41- ret = aclrtCreateStream(stream);42+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
42- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);43+ ret = aclrtCreateStream(stream);
43- return 0;44+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
45+ return 0;
44}46}
45 47 
46-template<typename T>48+template <typename T>
47int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,49int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr,
48- aclDataType dataType, aclTensor** tensor) {50+ aclDataType dataType, aclTensor** tensor)
49- auto size = GetShapeSize(shape) * sizeof(T);51+{
50- // 调用aclrtMalloc申请device侧引擎52+ auto size = GetShapeSize(shape) * sizeof(T);
51- auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);53+ // 调用aclrtMalloc申请device侧引擎
52- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);54+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
55+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret);
53 56 
54- // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上57+ // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上
55- ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);58+ ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE);
56- if (ret != ACL_SUCCESS) {59+ if (ret != ACL_SUCCESS) {
57- LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret);60+ LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret);
58- aclrtFree(*deviceAddr);61+ aclrtFree(*deviceAddr);
59- *deviceAddr = nullptr;62+ *deviceAddr = nullptr;
60- return ret;63+ return ret;
61- }64+ }
62 65 
63- // 计算连续tensor的strides66+ // 计算连续tensor的strides
64- std::vector<int64_t> strides(shape.size(), 1);67+ std::vector<int64_t> strides(shape.size(), 1);
65- for (int64_t i = shape.size() - 2; i >= 0; i--) {68+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
66- strides[i] = shape[i + 1] * strides[i + 1];69+ strides[i] = shape[i + 1] * strides[i + 1];
67- }70+ }
68 71 
69- // 调用aclCreateTensor接口创建aclTensor72+ // 调用aclCreateTensor接口创建aclTensor
70- *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,73+ *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND,
71- shape.data(), shape.size(), *deviceAddr);74+ shape.data(), shape.size(), *deviceAddr);
72- if (*tensor == nullptr) {75+ if (*tensor == nullptr) {
73- LOG_PRINT("aclCreateTensor failed.\n");76+ LOG_PRINT("aclCreateTensor failed.\n");
74- aclrtFree(*deviceAddr);77+ aclrtFree(*deviceAddr);
75- *deviceAddr = nullptr;78+ *deviceAddr = nullptr;
76- return -1;79+ return -1;
77- }80+ }
78- return 0;81+ return 0;
79}82}
80 83 
81-int main() {84+int main()
82- // 1. （固定写法）device/stream初始化, 参考acl API手册85+{
83- // 根据自己的实际device填写deviceId86+ // 1. （固定写法）device/stream初始化, 参考acl API手册
84- int32_t deviceId = 0;87+ // 根据自己的实际device填写deviceId
85- aclrtStream stream;88+ int32_t deviceId = 0;
86- auto ret = Init(deviceId, &stream);89+ aclrtStream stream;
87- CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);90+ auto ret = Init(deviceId, &stream);
88- // 2. 构造输入与输出，需要根据API的接口自定义构造91+ CHECK_RET(ret == 0, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret);
89- std::vector<int64_t> selfShape = {4, 2};92+ // 2. 构造输入与输出，需要根据API的接口自定义构造
90- std::vector<int64_t> gradOutputShape = {4, 2};93+ std::vector<int64_t> selfShape = {4, 2};
91- std::vector<int64_t> gradInputShape = {4, 2};94+ std::vector<int64_t> gradOutputShape = {4, 2};
92- void* selfDeviceAddr = nullptr;95+ std::vector<int64_t> gradInputShape = {4, 2};
93- void* gradOutputDeviceAddr = nullptr;96+ void* selfDeviceAddr = nullptr;
94- void* gradInputDeviceAddr = nullptr;97+ void* gradOutputDeviceAddr = nullptr;
95- aclTensor* self = nullptr;98+ void* gradInputDeviceAddr = nullptr;
96- aclTensor* gradOutput = nullptr;99+ aclTensor* self = nullptr;
97- aclTensor* gradInput = nullptr;100+ aclTensor* gradOutput = nullptr;
98- std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7};101+ aclTensor* gradInput = nullptr;
99- std::vector<float> gradOutputHostData = {1, 1, 1, 1, 1, 1, 1, 1};102+ std::vector<float> selfHostData = {0, 1, 2, 3, 4, 5, 6, 7};
100- std::vector<float> gradInputHostData = {0, 0, 0, 0, 0, 0, 0, 0};103+ std::vector<float> gradOutputHostData = {1, 1, 1, 1, 1, 1, 1, 1};
104+ std::vector<float> gradInputHostData = {0, 0, 0, 0, 0, 0, 0, 0};
101 105 
102- ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);106+ ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self);
103- CHECK_RET(ret == ACL_SUCCESS, return ret);107+ CHECK_RET(ret == ACL_SUCCESS, return ret);
104- ret = CreateAclTensor(gradOutputHostData, gradOutputShape, &gradOutputDeviceAddr, aclDataType::ACL_FLOAT, &gradOutput);108+ ret = CreateAclTensor(gradOutputHostData, gradOutputShape, &gradOutputDeviceAddr, aclDataType::ACL_FLOAT,
105- CHECK_RET(ret == ACL_SUCCESS, return ret);109+ &gradOutput);
106- ret = CreateAclTensor(gradInputHostData, gradInputShape, &gradInputDeviceAddr, aclDataType::ACL_FLOAT, &gradInput);110+ CHECK_RET(ret == ACL_SUCCESS, return ret);
107- CHECK_RET(ret == ACL_SUCCESS, return ret);111+ ret = CreateAclTensor(gradInputHostData, gradInputShape, &gradInputDeviceAddr, aclDataType::ACL_FLOAT, &gradInput);
112+ CHECK_RET(ret == ACL_SUCCESS, return ret);
108 113 
109- // 3. 调用CANN算子库API，需要修改为具体的API114+ // 3. 调用CANN算子库API，需要修改为具体的API
110- uint64_t workspaceSize = 0;115+ uint64_t workspaceSize = 0;
111- aclOpExecutor* executor;116+ aclOpExecutor* executor;
112- // 调用aclnnSeluBackward第一段接口117+ // 调用aclnnSeluBackward第一段接口
113- ret = aclnnSeluBackwardGetWorkspaceSize(gradOutput, self, gradInput, &workspaceSize, &executor);118+ ret = aclnnSeluBackwardGetWorkspaceSize(gradOutput, self, gradInput, &workspaceSize, &executor);
114- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSeluBackwardGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);119+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSeluBackwardGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);
115 120 
116- // 根据第一段接口计算出的workspaceSize申请device内存121+ // 根据第一段接口计算出的workspaceSize申请device内存
117- void* workspaceAddr = nullptr;122+ void* workspaceAddr = nullptr;
118- if (workspaceSize > 0) {123+ if (workspaceSize > 0) {
119- ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);124+ ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
120- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);125+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret);
121- }126+ }
122- // 调用aclnnSeluBackward第二段接口127+ // 调用aclnnSeluBackward第二段接口
123- ret = aclnnSeluBackward(workspaceAddr, workspaceSize, executor, stream);128+ ret = aclnnSeluBackward(workspaceAddr, workspaceSize, executor, stream);
124- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSeluBackward failed. ERROR: %d\n", ret); return ret);129+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSeluBackward failed. ERROR: %d\n", ret); return ret);
125- // 4. （固定写法）同步等待任务执行结束130+ // 4. （固定写法）同步等待任务执行结束
126- ret = aclrtSynchronizeStream(stream);131+ ret = aclrtSynchronizeStream(stream);
127- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);132+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);
128- // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改133+ // 5. 获取输出的值，将device侧内存上的结果拷贝至host侧，需要根据具体API的接口定义修改
129- auto size = GetShapeSize(gradInputShape);134+ auto size = GetShapeSize(gradInputShape);
130- std::vector<float> resultData(size, 0);135+ std::vector<float> resultData(size, 0);
131- ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), gradInputDeviceAddr, size * sizeof(float),136+ ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), gradInputDeviceAddr,
132- ACL_MEMCPY_DEVICE_TO_HOST);137+ size * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST);
133- CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);138+ CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret);
134- for (int64_t i = 0; i < size; i++) {139+ for (int64_t i = 0; i < size; i++) {
135- LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);140+ LOG_PRINT("result[%ld] is: %f\n", i, resultData[i]);
136- }141+ }
137- // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改142+ // 6. 释放aclTensor和aclScalar，需要根据具体API的接口定义修改
138- aclDestroyTensor(gradOutput);143+ aclDestroyTensor(gradOutput);
139- aclDestroyTensor(self);144+ aclDestroyTensor(self);
140- aclDestroyTensor(gradInput);145+ aclDestroyTensor(gradInput);
141 146 
142- // 7. 释放device资源，需要根据具体API的接口定义修改147+ // 7. 释放device资源，需要根据具体API的接口定义修改
143- aclrtFree(selfDeviceAddr);148+ aclrtFree(selfDeviceAddr);
144- aclrtFree(gradOutputDeviceAddr);149+ aclrtFree(gradOutputDeviceAddr);
145- aclrtFree(gradInputDeviceAddr);150+ aclrtFree(gradInputDeviceAddr);
146- if (workspaceSize > 0) {151+ if (workspaceSize > 0) {
147- aclrtFree(workspaceAddr);152+ aclrtFree(workspaceAddr);
148- }153+ }
149- aclrtDestroyStream(stream);154+ aclrtDestroyStream(stream);
150- aclrtResetDevice(deviceId);155+ aclrtResetDevice(deviceId);
151- aclFinalize();156+ aclFinalize();
152- return 0;157+ return 0;
153-}158+}
Mforeach/foreach_addcmul_scalar/op_kernel/foreach_addcmul_scalar.cpp+7-8
@@ -1,12 +1,11 @@
1/**1/**
2- * This program is free software, you can redistribute it and/or modify.
3 * Copyright (c) 2025 Huawei Technologies Co., Ltd.2 * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4- * This file is a part of the CANN Open Software.3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5- * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").4+ * 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, INCLUDING6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8- * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9- * 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.
10 */9 */
11 10 
12/*!11/*!
@@ -37,12 +36,12 @@ __aicore__ void AddcMulScalarAdapterForFloat(const LocalTensor<T>& tensorLocal1,
37template <typename T>36template <typename T>
38__aicore__ void AddcMulScalarAdapterForInt(const LocalTensor<T>& tensorLocal1, const LocalTensor<T>& tensorLocal2,37__aicore__ void AddcMulScalarAdapterForInt(const LocalTensor<T>& tensorLocal1, const LocalTensor<T>& tensorLocal2,
39 const LocalTensor<T>& tensorLocal3, const LocalTensor<float>& float32Tensor,38 const LocalTensor<T>& tensorLocal3, const LocalTensor<float>& float32Tensor,
40- const float scalarValue, const uint32_t maxCastDataCount,39+ const int32_t scalarValue, const uint32_t maxCastDataCount,
41 const int64_t dataCount)40 const int64_t dataCount)
42{41{
43 Mul(tensorLocal2, tensorLocal2, tensorLocal3, dataCount);42 Mul(tensorLocal2, tensorLocal2, tensorLocal3, dataCount);
44 PipeBarrier<PIPE_V>();43 PipeBarrier<PIPE_V>();
45- Muls(tensorLocal2, tensorLocal2, (int32_t)scalarValue, dataCount);44+ Muls(tensorLocal2, tensorLocal2, scalarValue, dataCount);
46 PipeBarrier<PIPE_V>();45 PipeBarrier<PIPE_V>();
47 Add(tensorLocal1, tensorLocal1, tensorLocal2, dataCount);46 Add(tensorLocal1, tensorLocal1, tensorLocal2, dataCount);
48}47}
Mforeach/foreach_addcmul_scalar_list/op_kernel/foreach_addcmul_scalar_list.cpp+3-3
@@ -38,12 +38,12 @@ __aicore__ void AddcMulScalarListAdapterForFloat(const LocalTensor<T>& tensor1Lo
38template <typename T>38template <typename T>
39__aicore__ void AddcMulScalarListAdapterForInt(const LocalTensor<T>& tensor1Local, const LocalTensor<T>& tensor2Local,39__aicore__ void AddcMulScalarListAdapterForInt(const LocalTensor<T>& tensor1Local, const LocalTensor<T>& tensor2Local,
40 const LocalTensor<T>& tensor3Local,40 const LocalTensor<T>& tensor3Local,
41- const LocalTensor<float>& float32Tensor, const float scalarValue,41+ const LocalTensor<float>& float32Tensor, const int32_t scalarValue,
42 const uint32_t maxCastDataCount, const int64_t dataCount)42 const uint32_t maxCastDataCount, const int64_t dataCount)
43{43{
44 Mul(tensor2Local, tensor2Local, tensor3Local, dataCount);44 Mul(tensor2Local, tensor2Local, tensor3Local, dataCount);
45 PipeBarrier<PIPE_V>();45 PipeBarrier<PIPE_V>();
46- Muls(tensor2Local, tensor2Local, (int32_t)scalarValue, dataCount);46+ Muls(tensor2Local, tensor2Local, scalarValue, dataCount);
47 PipeBarrier<PIPE_V>();47 PipeBarrier<PIPE_V>();
48 Add(tensor1Local, tensor1Local, tensor2Local, dataCount);48 Add(tensor1Local, tensor1Local, tensor2Local, dataCount);
49}49}
@@ -121,4 +121,4 @@ extern "C" __global__ __aicore__ void foreach_addcmul_scalar_list(GM_ADDR tensor
121#endif121#endif
122 }122 }
123#endif123#endif
124-}124+}
Mforeach/foreach_utils/op_kernel/foreach_one_scalar_list_quaternary_implict_output.h+4-3
@@ -44,7 +44,7 @@ protected:
44 GlobalTensor<T> inScalarGM;44 GlobalTensor<T> inScalarGM;
45 GM_ADDR inTensorsPtr_2 = nullptr;45 GM_ADDR inTensorsPtr_2 = nullptr;
46 GM_ADDR inTensorsPtr_3 = nullptr;46 GM_ADDR inTensorsPtr_3 = nullptr;
47- float scalarValue = 0.0;47+ ScalarT<T> scalarValue = 0;
48 48 
49private:49private:
50 __aicore__ inline void Compute(uint32_t index, int64_t dataCount, LocalTensor<float>& float32Tensor,50 __aicore__ inline void Compute(uint32_t index, int64_t dataCount, LocalTensor<float>& float32Tensor,
@@ -104,7 +104,8 @@ private:
104 104 
105 __aicore__ inline void valueScalar(const half& bVal) { scalarValue = (float)bVal; }105 __aicore__ inline void valueScalar(const half& bVal) { scalarValue = (float)bVal; }
106 106 
107- __aicore__ inline void valueScalar(const int& bVal) { scalarValue = static_cast<float>(bVal); }107+ // int32 标量按原始 int32_t 存储，避免 float() 强转丢精度 (issue #4519)
108+ __aicore__ inline void valueScalar(const int& bVal) { scalarValue = static_cast<int32_t>(bVal); }
108 109 
109 __aicore__ inline void valueScalar(const float& bVal) { scalarValue = bVal; }110 __aicore__ inline void valueScalar(const float& bVal) { scalarValue = bVal; }
110 111 
@@ -206,4 +207,4 @@ __aicore__ inline void ForeachOneScalarListQuaternaryImplictOutput<T, op, buffer
206} // namespace OpKernel207} // namespace OpKernel
207} // namespace Common208} // namespace Common
208 209 
209-#endif // KERNEL_FOREACH_ONE_SCALAR_LIST_QUATERNARY_IMPLICT_OUTPUT_H210+#endif // KERNEL_FOREACH_ONE_SCALAR_LIST_QUATERNARY_IMPLICT_OUTPUT_H
Mforeach/foreach_utils/op_kernel/foreach_one_scalar_quaternary_implict_output.h+24-5
@@ -22,15 +22,33 @@ namespace Common {
22namespace OpKernel {22namespace OpKernel {
23using namespace AscendC;23using namespace AscendC;
24 24 
25+// 标量存储类型萃取：int32 用原始 int32_t 存储避免 float() 强转丢精度
26+// half/bf16/float 仍用 float（计算本就提升到 float）。
27+template <typename T>
28+struct ScalarStorageType {
29+ using type = float;
30+};
31+template <>
32+struct ScalarStorageType<int32_t> {
33+ using type = int32_t;
34+};
35+// 注意：不再单独特化 ScalarStorageType<int>。
36+// 在该平台 int32_t 即为 int 的 typedef，同时特化二者会导致 "redefinition" 编译错误。
37+// int32_t 特化已覆盖 int。
38+ 
39+template <typename T>
40+using ScalarT = typename ScalarStorageType<T>::type;
41+ 
25template <typename T>42template <typename T>
26using OneScalarQuaternaryImplictOutputOp = void(const LocalTensor<T>&, const LocalTensor<T>&, const LocalTensor<T>&,43using OneScalarQuaternaryImplictOutputOp = void(const LocalTensor<T>&, const LocalTensor<T>&, const LocalTensor<T>&,
27- const LocalTensor<float>&, const float, const uint32_t, const int64_t);44+ const LocalTensor<float>&, const ScalarT<T>, const uint32_t,
45+ const int64_t);
28 46 
29template <typename T, OneScalarQuaternaryImplictOutputOp<T>* op>47template <typename T, OneScalarQuaternaryImplictOutputOp<T>* op>
30class InnerComputer {48class InnerComputer {
31public:49public:
32 __aicore__ inline void Compute(LocalTensor<T>& inLocal_1, LocalTensor<T>& inLocal_2, LocalTensor<T>& inLocal_3,50 __aicore__ inline void Compute(LocalTensor<T>& inLocal_1, LocalTensor<T>& inLocal_2, LocalTensor<T>& inLocal_3,
33- LocalTensor<float>& float32Tensor, float scalarVal, uint32_t maxCastDataCount,51+ LocalTensor<float>& float32Tensor, ScalarT<T> scalarVal, uint32_t maxCastDataCount,
34 int64_t dataCount)52 int64_t dataCount)
35 {53 {
36 PipeBarrier<PIPE_V>();54 PipeBarrier<PIPE_V>();
@@ -61,7 +79,7 @@ protected:
61 GlobalTensor<DTYPE_SCALAR> inScalarGM;79 GlobalTensor<DTYPE_SCALAR> inScalarGM;
62 GM_ADDR inTensorsPtr_2 = nullptr;80 GM_ADDR inTensorsPtr_2 = nullptr;
63 GM_ADDR inTensorsPtr_3 = nullptr;81 GM_ADDR inTensorsPtr_3 = nullptr;
64- float scalarVal = 0.0;82+ ScalarT<T> scalarVal = 0;
65 83 
66private:84private:
67 __aicore__ inline void Compute(uint32_t index, int64_t dataCount, LocalTensor<float>& float32Tensor,85 __aicore__ inline void Compute(uint32_t index, int64_t dataCount, LocalTensor<float>& float32Tensor,
@@ -132,7 +150,8 @@ __aicore__ inline void ForeachOneScalarQuaternaryImplictOutput<T, op, bufferNum,
132 inTensorsPtr_3 = x3;150 inTensorsPtr_3 = x3;
133 Base::outTensorsPtr = y;151 Base::outTensorsPtr = y;
134 inScalarGM.SetGlobalBuffer((__gm__ DTYPE_SCALAR*)scalar, 1);152 inScalarGM.SetGlobalBuffer((__gm__ DTYPE_SCALAR*)scalar, 1);
135- scalarVal = float(inScalarGM.GetValue(0));153+ // int32 按 int32_t 存储避免 float() 强转丢精度；half/bf16/float 经 float 中转
154+ scalarVal = ScalarT<T>(inScalarGM.GetValue(0));
136#if __CCE_AICORE__ >= 220155#if __CCE_AICORE__ >= 220
137 if (std::is_same_v<T, bfloat16_t> || std::is_same_v<T, half>) {156 if (std::is_same_v<T, bfloat16_t> || std::is_same_v<T, half>) {
138 Base::Base::pipe.InitBuffer(Base::float32Queue, 1, Base::Base::inputsTensorUbSize * paramsCount);157 Base::Base::pipe.InitBuffer(Base::float32Queue, 1, Base::Base::inputsTensorUbSize * paramsCount);
@@ -214,4 +233,4 @@ __aicore__ inline void ForeachOneScalarQuaternaryImplictOutput<T, op, bufferNum,
214} // namespace OpKernel233} // namespace OpKernel
215} // namespace Common234} // namespace Common
216 235 
217-#endif // KERNEL_FOREACH_ONE_SCALAR_QUATERNARY_IMPLICT_OUTPUT_H236+#endif // KERNEL_FOREACH_ONE_SCALAR_QUATERNARY_IMPLICT_OUTPUT_H