Star1738
1648
Code
Code
Issues469
Pull Requests550
Actions
Discussion
Wiki
Analysis
Star1738
1648
Merged
feat(random): stateless_random_normal_v2/v3、stateless_uniform 支持 ascend960dt/ascend960pr 平台 #6175
吴成文Add on Oct 9
feat(random): stateless_random_normal_v2/v3、stateless_uniform 支持 ascend960dt/ascend960pr 平台 #6175
Merged
吴成文Add on Oct 9
Merged fromwuchengwen_oh/ops-mathfeat/random-960-supportintocann/ops-mathmaster
Total: 24 files changed+1895-64
Mrandom/stateless_random_normal_v2/CMakeLists.txt+2-2
@@ -9,7 +9,7 @@
9# ---------------------------------------------------------------------------------------------------------9# ---------------------------------------------------------------------------------------------------------
10 10 
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12-set(SUPPORT_COMPUTE_UNIT "ascend950" "ascend350")12+set(SUPPORT_COMPUTE_UNIT "ascend950" "ascend350" "ascend960dt" "ascend960pr")
G
GgubaochengOct 9

ascend960dt/960pr 未登记全局构建表。本 PR 已在三算子 CMakeLists、def.cpp、ascendc_config.json 加 960，但根 CMakeLists.txt:77 ASCEND_ALL_COMPUTE_UNIT 与 scripts/build.conf.sh:14 SUPPORT_COMPUTE_UNIT_SHORT 均不含这两颗芯片。后果同 #6174：gen_ops_info.cmake 无法生成 960 注册清单，build_cmake.sh 校验 --soc= 时 exit 1，标准构建路径无法产出 960 二进制。

（来自自动检视代理，以 gubaocheng 身份提交）

likedislike
13# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch35" "arch35" "arch35" "arch35")
15add_all_modules_sources(OPTYPE stateless_random_normal_v2 ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DEPENDENCIES random_common DISABLE_IN_OPP TRUE)15add_all_modules_sources(OPTYPE stateless_random_normal_v2 ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DEPENDENCIES random_common DISABLE_IN_OPP TRUE)
Mrandom/stateless_random_normal_v2/op_api/aclnn_normal_out.cpp+8-4
@@ -43,6 +43,12 @@ static const std::initializer_list<op::DataType> DTYPE_SUPPORT_LIST = {op::DataT
43// 表示mean,std的数据类型是Tensor还是Scalar43// 表示mean,std的数据类型是Tensor还是Scalar
44enum class ScalarMode { TensorTensor, TensorScalar, ScalarTensor, ScalarScalar };44enum class ScalarMode { TensorTensor, TensorScalar, ScalarTensor, ScalarScalar };
45 45 
46+static inline bool IsNormalV3SupportedArch()
47+{
48+ auto curArch = GetCurrentPlatformInfo().GetCurNpuArch();
49+ return curArch == NpuArch::DAV_3510 || curArch == NpuArch::DAV_9201 || curArch == NpuArch::DAV_9202;
50+}
51+ 
46/* 查看TensorFloat的Dtype和Shape */52/* 查看TensorFloat的Dtype和Shape */
47static bool CheckTensorAndFloatDtype(const aclTensor* mean, const aclTensor* out)53static bool CheckTensorAndFloatDtype(const aclTensor* mean, const aclTensor* out)
48{54{
@@ -211,8 +217,7 @@ aclnnStatus CommonLogicGeneralNormal(const aclTensor* mean, const aclTensor* std
211{217{
212 const aclTensor* addOut = nullptr;218 const aclTensor* addOut = nullptr;
213 219 
214- if (!aclnnGetPytorchRandom() && GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_3510 &&220+ if (!aclnnGetPytorchRandom() && IsNormalV3SupportedArch() && self->GetDataType() != DataType::DT_DOUBLE) {
215- self->GetDataType() != DataType::DT_DOUBLE) {
216 OP_LOGD("compat mode, use V3 Normal");221 OP_LOGD("compat mode, use V3 Normal");
217 // seed转化为key222 // seed转化为key
218 FVector<int64_t, op::MAX_DIM_NUM> key_vec = {seed};223 FVector<int64_t, op::MAX_DIM_NUM> key_vec = {seed};
@@ -253,8 +258,7 @@ aclnnStatus CommonLogicGeneralNormal(const aclTensor* mean, const aclTensor* std
253 }258 }
254 }259 }
255 addOut = l0op::StatelessRandomNormalV3(self, keyArr, counterArr, meanCasted, stdCasted, uniqueExecutor.get());260 addOut = l0op::StatelessRandomNormalV3(self, keyArr, counterArr, meanCasted, stdCasted, uniqueExecutor.get());
256- } else if (GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_3510 &&261+ } else if (IsNormalV3SupportedArch() && self->GetDataType() != DataType::DT_DOUBLE) {
257- self->GetDataType() != DataType::DT_DOUBLE) {
258 if (scalarMode == ScalarMode::TensorTensor || scalarMode == ScalarMode::ScalarTensor) {262 if (scalarMode == ScalarMode::TensorTensor || scalarMode == ScalarMode::ScalarTensor) {
259 FVector<float> zeroMeanVector = {0.0f};263 FVector<float> zeroMeanVector = {0.0f};
260 auto zeroMeanTensor = uniqueExecutor.get()->ConvertToTensor(zeroMeanVector.data(), zeroMeanVector.size(),264 auto zeroMeanTensor = uniqueExecutor.get()->ConvertToTensor(zeroMeanVector.data(), zeroMeanVector.size(),
Mrandom/stateless_random_normal_v2/op_host/arch35/stateless_random_normal_v2_tiling_arch35.cpp+13-3
@@ -15,6 +15,7 @@
15#include "stateless_random_normal_v2_tiling_arch35.h"15#include "stateless_random_normal_v2_tiling_arch35.h"
16#include <string>16#include <string>
17#include "platform/platform_info.h"17#include "platform/platform_info.h"
18+#include "platform/platform_ascendc.h"
18#include "log/log.h"19#include "log/log.h"
19#include "register/op_impl_registry.h"20#include "register/op_impl_registry.h"
20 21 
@@ -32,6 +33,8 @@ static constexpr uint16_t SIZE_OF_FLOAT = 4;
32static constexpr uint16_t SPLIT_UB_NUM = 5;33static constexpr uint16_t SPLIT_UB_NUM = 5;
33static constexpr int64_t COUNTER_NUMBER_LOW_BOUND = 1;34static constexpr int64_t COUNTER_NUMBER_LOW_BOUND = 1;
34 35 
36+static inline bool IsArch92(NpuArch npuArch) { return npuArch == NpuArch::DAV_9201 || npuArch == NpuArch::DAV_9202; }
37+ 
35ge::graphStatus StatelessRandomNormalV2Tiling::GetPlatformInfo()38ge::graphStatus StatelessRandomNormalV2Tiling::GetPlatformInfo()
36{39{
37 auto platformInfo = context_->GetPlatformInfo();40 auto platformInfo = context_->GetPlatformInfo();
@@ -40,6 +43,7 @@ ge::graphStatus StatelessRandomNormalV2Tiling::GetPlatformInfo()
40 OP_CHECK_IF(compileInfoPtr == nullptr, OP_LOGE(context_, "compile info is null"), return ge::GRAPH_FAILED);43 OP_CHECK_IF(compileInfoPtr == nullptr, OP_LOGE(context_, "compile info is null"), return ge::GRAPH_FAILED);
41 coreNum_ = compileInfoPtr->aivNum;44 coreNum_ = compileInfoPtr->aivNum;
42 ubSize_ = compileInfoPtr->ubSize - REGBASE_CCEC_CACHE_SIZE;45 ubSize_ = compileInfoPtr->ubSize - REGBASE_CCEC_CACHE_SIZE;
46+ isArch92_ = compileInfoPtr->isArch92;
43 } else {47 } else {
44 auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);48 auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
45 auto aivNum = ascendcPlatform.GetCoreNumAiv();49 auto aivNum = ascendcPlatform.GetCoreNumAiv();
@@ -53,6 +57,7 @@ ge::graphStatus StatelessRandomNormalV2Tiling::GetPlatformInfo()
53 REGBASE_CCEC_CACHE_SIZE),57 REGBASE_CCEC_CACHE_SIZE),
54 return ge::GRAPH_FAILED);58 return ge::GRAPH_FAILED);
55 ubSize_ = ubSizePlatForm - REGBASE_CCEC_CACHE_SIZE;59 ubSize_ = ubSizePlatForm - REGBASE_CCEC_CACHE_SIZE;
60+ isArch92_ = IsArch92(ascendcPlatform.GetCurNpuArch());
56 }61 }
57 return ge::GRAPH_SUCCESS;62 return ge::GRAPH_SUCCESS;
58}63}
@@ -133,7 +138,7 @@ inline T1 StatelessRandomNormalV2Tiling::CeilDiv(const T1 a, const T2 b) const
133void StatelessRandomNormalV2Tiling::BlockTiling()138void StatelessRandomNormalV2Tiling::BlockTiling()
134{139{
135 outputDtypeSize_ = SIZE_OF_FLOAT;140 outputDtypeSize_ = SIZE_OF_FLOAT;
136- auto coreAlignFactor = CORE_ALIGN_SIZE / outputDtypeSize_;141+ auto coreAlignFactor = isArch92_ ? CORE_ALIGN_ELEMENTS_ARCH92 : CORE_ALIGN_SIZE / outputDtypeSize_;
137 auto blockFactor = CeilDiv(outputSize_, coreNum_);142 auto blockFactor = CeilDiv(outputSize_, coreNum_);
138 auto blockAlignFactor = CeilDiv(blockFactor, coreAlignFactor) * coreAlignFactor;143 auto blockAlignFactor = CeilDiv(blockFactor, coreAlignFactor) * coreAlignFactor;
139 auto minTilingSize = MIN_TILING_SIZE;144 auto minTilingSize = MIN_TILING_SIZE;
@@ -151,8 +156,12 @@ ge::graphStatus StatelessRandomNormalV2Tiling::UbTiling()
151{156{
152 // splitUbSize: 2 for double buffer; 3 for data converse157 // splitUbSize: 2 for double buffer; 3 for data converse
153 auto splitUbSize = ubSize_ / SPLIT_UB_NUM;158 auto splitUbSize = ubSize_ / SPLIT_UB_NUM;
154- auto alignFactor = BLOCK_SIZE_BYTES / outputDtypeSize_;159+ if (isArch92_) {
155- ubTilingSize_ = CeilDiv(splitUbSize / outputDtypeSize_, alignFactor) * alignFactor;160+ ubTilingSize_ = splitUbSize / outputDtypeSize_ / UB_ALIGN_ELEMENTS_ARCH92 * UB_ALIGN_ELEMENTS_ARCH92;
161+ } else {
162+ auto alignFactor = BLOCK_SIZE_BYTES / outputDtypeSize_;
163+ ubTilingSize_ = CeilDiv(splitUbSize / outputDtypeSize_, alignFactor) * alignFactor;
164+ }
156 OP_LOGD(opName, "splitUbSize = %u, ubTilingSize = %u", splitUbSize, ubTilingSize_);165 OP_LOGD(opName, "splitUbSize = %u, ubTilingSize = %u", splitUbSize, ubTilingSize_);
157 return ge::GRAPH_SUCCESS;166 return ge::GRAPH_SUCCESS;
158}167}
@@ -227,6 +236,7 @@ static ge::graphStatus TilingPrepare4StatelessRandomNormalV2(gert::TilingParseCo
227 auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);236 auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
228 compileInfoPtr->aivNum = ascendcPlatform.GetCoreNumAiv();237 compileInfoPtr->aivNum = ascendcPlatform.GetCoreNumAiv();
229 ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfoPtr->ubSize);238 ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfoPtr->ubSize);
239+ compileInfoPtr->isArch92 = IsArch92(ascendcPlatform.GetCurNpuArch());
230 return ge::GRAPH_SUCCESS;240 return ge::GRAPH_SUCCESS;
231}241}
232 242 
Mrandom/stateless_random_normal_v2/op_host/arch35/stateless_random_normal_v2_tiling_arch35.h+8-13
@@ -38,25 +38,17 @@ REGISTER_TILING_DATA_CLASS(StatelessRandomNormalV2, StatelessRandomNormalV2Tilin
38struct StatelessRandomNormalV2CompileInfo {38struct StatelessRandomNormalV2CompileInfo {
39 uint64_t aivNum;39 uint64_t aivNum;
40 uint64_t ubSize;40 uint64_t ubSize;
41+ bool isArch92 = false;
41};42};
42 43 
43class StatelessRandomNormalV2Tiling : public Ops::Base::TilingBaseClass {44class StatelessRandomNormalV2Tiling : public Ops::Base::TilingBaseClass {
44public:45public:
45- explicit StatelessRandomNormalV2Tiling(gert::TilingContext* context) : Ops::Base::TilingBaseClass(context)46+ explicit StatelessRandomNormalV2Tiling(gert::TilingContext* context) : Ops::Base::TilingBaseClass(context) {}
46- {}
47 47 
48- enum class Algorithm : int32_t48+ enum class Algorithm : int32_t { RNG_ALG_PHILOX = 1, RNG_ALG_THREEFRY = 2, RNG_ALG_AUTO_SELECT = 3 };
49- {
50- RNG_ALG_PHILOX = 1,
51- RNG_ALG_THREEFRY = 2,
52- RNG_ALG_AUTO_SELECT = 3
53- };
54 49 
55protected:50protected:
56- bool IsCapable() override51+ bool IsCapable() override { return true; }
57- {
58- return true;
59- }
60 52 
61 // 1、获取平台信息比如CoreNum、UB/L1/L0C资源大小53 // 1、获取平台信息比如CoreNum、UB/L1/L0C资源大小
62 ge::graphStatus GetPlatformInfo() override;54 ge::graphStatus GetPlatformInfo() override;
@@ -76,6 +68,8 @@ protected:
76private:68private:
77 static constexpr uint32_t CORE_ALIGN_SIZE = 512;69 static constexpr uint32_t CORE_ALIGN_SIZE = 512;
78 static constexpr uint32_t BLOCK_SIZE_BYTES = 32;70 static constexpr uint32_t BLOCK_SIZE_BYTES = 32;
71+ static constexpr uint32_t CORE_ALIGN_ELEMENTS_ARCH92 = 32;
72+ static constexpr uint32_t UB_ALIGN_ELEMENTS_ARCH92 = 256;
79 static constexpr uint32_t MIN_TILING_SIZE = 256;73 static constexpr uint32_t MIN_TILING_SIZE = 256;
80 static constexpr uint32_t REGBASE_CCEC_CACHE_SIZE = 8 * 1024;74 static constexpr uint32_t REGBASE_CCEC_CACHE_SIZE = 8 * 1024;
81 static constexpr uint32_t DEFAULT_WORKSPACE_SIZE = 16 * 1024 * 1024;75 static constexpr uint32_t DEFAULT_WORKSPACE_SIZE = 16 * 1024 * 1024;
@@ -91,6 +85,7 @@ private:
91private:85private:
92 uint32_t coreNum_ = 0;86 uint32_t coreNum_ = 0;
93 uint32_t ubSize_ = 0;87 uint32_t ubSize_ = 0;
88+ bool isArch92_ = false;
94 ge::DataType outputDtype_ = ge::DT_UNDEFINED;89 ge::DataType outputDtype_ = ge::DT_UNDEFINED;
95 uint32_t outputDtypeVal_ = 0;90 uint32_t outputDtypeVal_ = 0;
96 uint64_t outputSize_ = 1;91 uint64_t outputSize_ = 1;
@@ -106,4 +101,4 @@ private:
106 StatelessRandomNormalV2TilingData tilingData;101 StatelessRandomNormalV2TilingData tilingData;
107};102};
108 103 
109-} // namespace optiling104+} // namespace optiling
Arandom/stateless_random_normal_v2/op_host/config/ascend960dt/stateless_random_normal_v2_binary.json+389-0
@@ -0,0 +1,389 @@
1+{
2+ "op_type": "StatelessRandomNormalV2",
3+ "op_list": [
4+ {
5+ "bin_filename": "StatelessRandomNormalV2_1ea4bdp73cd9df67072d9z7aaec9c583",
6+ "inputs": [
7+ {
8+ "name": "shape",
9+ "index": 0,
10+ "dtype": "int32",
11+ "format": "ND",
12+ "shape": [
13+ -2
14+ ],
15+ "paramType": "required"
16+ },
17+ {
18+ "name": "key",
19+ "index": 1,
20+ "dtype": "uint64",
21+ "format": "ND",
22+ "shape": [
23+ -2
24+ ],
25+ "paramType": "required"
26+ },
27+ {
28+ "name": "counter",
29+ "index": 2,
30+ "dtype": "uint64",
31+ "format": "ND",
32+ "shape": [
33+ -2
34+ ],
35+ "paramType": "required"
36+ },
37+ {
38+ "name": "alg",
39+ "index": 3,
40+ "dtype": "int32",
41+ "format": "ND",
42+ "shape": [
43+ -2
44+ ],
45+ "paramType": "required"
46+ }
47+ ],
48+ "outputs": [
49+ {
50+ "name": "y",
51+ "index": 0,
52+ "dtype": "float32",
53+ "format": "ND",
54+ "shape": [
55+ -2
56+ ],
57+ "paramType": "required"
58+ }
59+ ],
60+ "attrs": [
61+ {
62+ "name": "dtype",
63+ "dtype": "int",
64+ "value": null
65+ }
66+ ]
67+ },
68+ {
69+ "bin_filename": "StatelessRandomNormalV2_1ea4bd0e3cd9df67072d9z7akyc9c583",
70+ "inputs": [
71+ {
72+ "name": "shape",
73+ "index": 0,
74+ "dtype": "int32",
75+ "format": "ND",
76+ "shape": [
77+ -2
78+ ],
79+ "paramType": "required"
80+ },
81+ {
82+ "name": "key",
83+ "index": 1,
84+ "dtype": "uint64",
85+ "format": "ND",
86+ "shape": [
87+ -2
88+ ],
89+ "paramType": "required"
90+ },
91+ {
92+ "name": "counter",
93+ "index": 2,
94+ "dtype": "uint64",
95+ "format": "ND",
96+ "shape": [
97+ -2
98+ ],
99+ "paramType": "required"
100+ },
101+ {
102+ "name": "alg",
103+ "index": 3,
104+ "dtype": "int32",
105+ "format": "ND",
106+ "shape": [
107+ -2
108+ ],
109+ "paramType": "required"
110+ }
111+ ],
112+ "outputs": [
113+ {
114+ "name": "y",
115+ "index": 0,
116+ "dtype": "float16",
117+ "format": "ND",
118+ "shape": [
119+ -2
120+ ],
121+ "paramType": "required"
122+ }
123+ ],
124+ "attrs": [
125+ {
126+ "name": "dtype",
127+ "dtype": "int",
128+ "value": null
129+ }
130+ ]
131+ },
132+ {
133+ "bin_filename": "StatelessRandomNormalV2_1ea4bd9g3cd9df67072d9z7amgc9c733",
134+ "inputs": [
135+ {
136+ "name": "shape",
137+ "index": 0,
138+ "dtype": "int32",
139+ "format": "ND",
140+ "shape": [
141+ -2
142+ ],
143+ "paramType": "required"
144+ },
145+ {
146+ "name": "key",
147+ "index": 1,
148+ "dtype": "uint64",
149+ "format": "ND",
150+ "shape": [
151+ -2
152+ ],
153+ "paramType": "required"
154+ },
155+ {
156+ "name": "counter",
157+ "index": 2,
158+ "dtype": "uint64",
159+ "format": "ND",
160+ "shape": [
161+ -2
162+ ],
163+ "paramType": "required"
164+ },
165+ {
166+ "name": "alg",
167+ "index": 3,
168+ "dtype": "int32",
169+ "format": "ND",
170+ "shape": [
171+ -2
172+ ],
173+ "paramType": "required"
174+ }
175+ ],
176+ "outputs": [
177+ {
178+ "name": "y",
179+ "index": 0,
180+ "dtype": "bfloat16",
181+ "format": "ND",
182+ "shape": [
183+ -2
184+ ],
185+ "paramType": "required"
186+ }
187+ ],
188+ "attrs": [
189+ {
190+ "name": "dtype",
191+ "dtype": "int",
192+ "value": null
193+ }
194+ ]
195+ },
196+ {
197+ "bin_filename": "StatelessRandomNormalV2_1ea4bd0e3cd9df67072d9z7aaec9c583",
198+ "inputs": [
199+ {
200+ "name": "shape",
201+ "index": 0,
202+ "dtype": "int64",
203+ "format": "ND",
204+ "shape": [
205+ -2
206+ ],
207+ "paramType": "required"
208+ },
209+ {
210+ "name": "key",
211+ "index": 1,
212+ "dtype": "uint64",
213+ "format": "ND",
214+ "shape": [
215+ -2
216+ ],
217+ "paramType": "required"
218+ },
219+ {
220+ "name": "counter",
221+ "index": 2,
222+ "dtype": "uint64",
223+ "format": "ND",
224+ "shape": [
225+ -2
226+ ],
227+ "paramType": "required"
228+ },
229+ {
230+ "name": "alg",
231+ "index": 3,
232+ "dtype": "int32",
233+ "format": "ND",
234+ "shape": [
235+ -2
236+ ],
237+ "paramType": "required"
238+ }
239+ ],
240+ "outputs": [
241+ {
242+ "name": "y",
243+ "index": 0,
244+ "dtype": "float32",
245+ "format": "ND",
246+ "shape": [
247+ -2
248+ ],
249+ "paramType": "required"
250+ }
251+ ],
252+ "attrs": [
253+ {
254+ "name": "dtype",
255+ "dtype": "int",
256+ "value": null
257+ }
258+ ]
259+ },
260+ {
261+ "bin_filename": "StatelessRandomNormalV2_1ea4bd0e3cd9df67072d9z7amgc9c583",
262+ "inputs": [
263+ {
264+ "name": "shape",
265+ "index": 0,
266+ "dtype": "int64",
267+ "format": "ND",
268+ "shape": [
269+ -2
270+ ],
271+ "paramType": "required"
272+ },
273+ {
274+ "name": "key",
275+ "index": 1,
276+ "dtype": "uint64",
277+ "format": "ND",
278+ "shape": [
279+ -2
280+ ],
281+ "paramType": "required"
282+ },
283+ {
284+ "name": "counter",
285+ "index": 2,
286+ "dtype": "uint64",
287+ "format": "ND",
288+ "shape": [
289+ -2
290+ ],
291+ "paramType": "required"
292+ },
293+ {
294+ "name": "alg",
295+ "index": 3,
296+ "dtype": "int32",
297+ "format": "ND",
298+ "shape": [
299+ -2
300+ ],
301+ "paramType": "required"
302+ }
303+ ],
304+ "outputs": [
305+ {
306+ "name": "y",
307+ "index": 0,
308+ "dtype": "float16",
309+ "format": "ND",
310+ "shape": [
311+ -2
312+ ],
313+ "paramType": "required"
314+ }
315+ ],
316+ "attrs": [
317+ {
318+ "name": "dtype",
319+ "dtype": "int",
320+ "value": null
321+ }
322+ ]
323+ },
324+ {
325+ "bin_filename": "StatelessRandomNormalV2_1ea4bd0e3cd9df67072d9z7amgc9c733",
326+ "inputs": [
327+ {
328+ "name": "shape",
329+ "index": 0,
330+ "dtype": "int64",
331+ "format": "ND",
332+ "shape": [
333+ -2
334+ ],
335+ "paramType": "required"
336+ },
337+ {
338+ "name": "key",
339+ "index": 1,
340+ "dtype": "uint64",
341+ "format": "ND",
342+ "shape": [
343+ -2
344+ ],
345+ "paramType": "required"
346+ },
347+ {
348+ "name": "counter",
349+ "index": 2,
350+ "dtype": "uint64",
351+ "format": "ND",
352+ "shape": [
353+ -2
354+ ],
355+ "paramType": "required"
356+ },
357+ {
358+ "name": "alg",
359+ "index": 3,
360+ "dtype": "int32",
361+ "format": "ND",
362+ "shape": [
363+ -2
364+ ],
365+ "paramType": "required"
366+ }
367+ ],
368+ "outputs": [
369+ {
370+ "name": "y",
371+ "index": 0,
372+ "dtype": "bfloat16",
373+ "format": "ND",
374+ "shape": [
375+ -2
376+ ],
377+ "paramType": "required"
378+ }
379+ ],
380+ "attrs": [
381+ {
382+ "name": "dtype",
383+ "dtype": "int",
384+ "value": null
385+ }
386+ ]
387+ }
388+ ]
389+}
Arandom/stateless_random_normal_v2/op_host/config/ascend960dt/stateless_random_normal_v2_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，ascnedxx为可选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+[StatelessRandomNormalV2]
13+default=0
Arandom/stateless_random_normal_v2/op_host/config/ascend960pr/stateless_random_normal_v2_binary.json+389-0
@@ -0,0 +1,389 @@
1+{
2+ "op_type": "StatelessRandomNormalV2",
3+ "op_list": [
4+ {
5+ "bin_filename": "StatelessRandomNormalV2_1ea4bdp73cd9df67072d9z7aaec9c583",
6+ "inputs": [
7+ {
8+ "name": "shape",
9+ "index": 0,
10+ "dtype": "int32",
11+ "format": "ND",
12+ "shape": [
13+ -2
14+ ],
15+ "paramType": "required"
16+ },
17+ {
18+ "name": "key",
19+ "index": 1,
20+ "dtype": "uint64",
21+ "format": "ND",
22+ "shape": [
23+ -2
24+ ],
25+ "paramType": "required"
26+ },
27+ {
28+ "name": "counter",
29+ "index": 2,
30+ "dtype": "uint64",
31+ "format": "ND",
32+ "shape": [
33+ -2
34+ ],
35+ "paramType": "required"
36+ },
37+ {
38+ "name": "alg",
39+ "index": 3,
40+ "dtype": "int32",
41+ "format": "ND",
42+ "shape": [
43+ -2
44+ ],
45+ "paramType": "required"
46+ }
47+ ],
48+ "outputs": [
49+ {
50+ "name": "y",
51+ "index": 0,
52+ "dtype": "float32",
53+ "format": "ND",
54+ "shape": [
55+ -2
56+ ],
57+ "paramType": "required"
58+ }
59+ ],
60+ "attrs": [
61+ {
62+ "name": "dtype",
63+ "dtype": "int",
64+ "value": null
65+ }
66+ ]
67+ },
68+ {
69+ "bin_filename": "StatelessRandomNormalV2_1ea4bd0e3cd9df67072d9z7akyc9c583",
70+ "inputs": [
71+ {
72+ "name": "shape",
73+ "index": 0,
74+ "dtype": "int32",
75+ "format": "ND",
76+ "shape": [
77+ -2
78+ ],
79+ "paramType": "required"
80+ },
81+ {
82+ "name": "key",
83+ "index": 1,
84+ "dtype": "uint64",
85+ "format": "ND",
86+ "shape": [
87+ -2
88+ ],
89+ "paramType": "required"
90+ },
91+ {
92+ "name": "counter",
93+ "index": 2,
94+ "dtype": "uint64",
95+ "format": "ND",
96+ "shape": [
97+ -2
98+ ],
99+ "paramType": "required"
100+ },
101+ {
102+ "name": "alg",
103+ "index": 3,
104+ "dtype": "int32",
105+ "format": "ND",
106+ "shape": [
107+ -2
108+ ],
109+ "paramType": "required"
110+ }
111+ ],
112+ "outputs": [
113+ {
114+ "name": "y",
115+ "index": 0,
116+ "dtype": "float16",
117+ "format": "ND",
118+ "shape": [
119+ -2
120+ ],
121+ "paramType": "required"
122+ }
123+ ],
124+ "attrs": [
125+ {
126+ "name": "dtype",
127+ "dtype": "int",
128+ "value": null
129+ }
130+ ]
131+ },
132+ {
133+ "bin_filename": "StatelessRandomNormalV2_1ea4bd9g3cd9df67072d9z7amgc9c733",
134+ "inputs": [
135+ {
136+ "name": "shape",
137+ "index": 0,
138+ "dtype": "int32",
139+ "format": "ND",
140+ "shape": [
141+ -2
142+ ],
143+ "paramType": "required"
144+ },
145+ {
146+ "name": "key",
147+ "index": 1,
148+ "dtype": "uint64",
149+ "format": "ND",
150+ "shape": [
151+ -2
152+ ],
153+ "paramType": "required"
154+ },
155+ {
156+ "name": "counter",
157+ "index": 2,
158+ "dtype": "uint64",
159+ "format": "ND",
160+ "shape": [
161+ -2
162+ ],
163+ "paramType": "required"
164+ },
165+ {
166+ "name": "alg",
167+ "index": 3,
168+ "dtype": "int32",
169+ "format": "ND",
170+ "shape": [
171+ -2
172+ ],
173+ "paramType": "required"
174+ }
175+ ],
176+ "outputs": [
177+ {
178+ "name": "y",
179+ "index": 0,
180+ "dtype": "bfloat16",
181+ "format": "ND",
182+ "shape": [
183+ -2
184+ ],
185+ "paramType": "required"
186+ }
187+ ],
188+ "attrs": [
189+ {
190+ "name": "dtype",
191+ "dtype": "int",
192+ "value": null
193+ }
194+ ]
195+ },
196+ {
197+ "bin_filename": "StatelessRandomNormalV2_1ea4bd0e3cd9df67072d9z7aaec9c583",
198+ "inputs": [
199+ {
200+ "name": "shape",
201+ "index": 0,
202+ "dtype": "int64",
203+ "format": "ND",
204+ "shape": [
205+ -2
206+ ],
207+ "paramType": "required"
208+ },
209+ {
210+ "name": "key",
211+ "index": 1,
212+ "dtype": "uint64",
213+ "format": "ND",
214+ "shape": [
215+ -2
216+ ],
217+ "paramType": "required"
218+ },
219+ {
220+ "name": "counter",
221+ "index": 2,
222+ "dtype": "uint64",
223+ "format": "ND",
224+ "shape": [
225+ -2
226+ ],
227+ "paramType": "required"
228+ },
229+ {
230+ "name": "alg",
231+ "index": 3,
232+ "dtype": "int32",
233+ "format": "ND",
234+ "shape": [
235+ -2
236+ ],
237+ "paramType": "required"
238+ }
239+ ],
240+ "outputs": [
241+ {
242+ "name": "y",
243+ "index": 0,
244+ "dtype": "float32",
245+ "format": "ND",
246+ "shape": [
247+ -2
248+ ],
249+ "paramType": "required"
250+ }
251+ ],
252+ "attrs": [
253+ {
254+ "name": "dtype",
255+ "dtype": "int",
256+ "value": null
257+ }
258+ ]
259+ },
260+ {
261+ "bin_filename": "StatelessRandomNormalV2_1ea4bd0e3cd9df67072d9z7amgc9c583",
262+ "inputs": [
263+ {
264+ "name": "shape",
265+ "index": 0,
266+ "dtype": "int64",
267+ "format": "ND",
268+ "shape": [
269+ -2
270+ ],
271+ "paramType": "required"
272+ },
273+ {
274+ "name": "key",
275+ "index": 1,
276+ "dtype": "uint64",
277+ "format": "ND",
278+ "shape": [
279+ -2
280+ ],
281+ "paramType": "required"
282+ },
283+ {
284+ "name": "counter",
285+ "index": 2,
286+ "dtype": "uint64",
287+ "format": "ND",
288+ "shape": [
289+ -2
290+ ],
291+ "paramType": "required"
292+ },
293+ {
294+ "name": "alg",
295+ "index": 3,
296+ "dtype": "int32",
297+ "format": "ND",
298+ "shape": [
299+ -2
300+ ],
301+ "paramType": "required"
302+ }
303+ ],
304+ "outputs": [
305+ {
306+ "name": "y",
307+ "index": 0,
308+ "dtype": "float16",
309+ "format": "ND",
310+ "shape": [
311+ -2
312+ ],
313+ "paramType": "required"
314+ }
315+ ],
316+ "attrs": [
317+ {
318+ "name": "dtype",
319+ "dtype": "int",
320+ "value": null
321+ }
322+ ]
323+ },
324+ {
325+ "bin_filename": "StatelessRandomNormalV2_1ea4bd0e3cd9df67072d9z7amgc9c733",
326+ "inputs": [
327+ {
328+ "name": "shape",
329+ "index": 0,
330+ "dtype": "int64",
331+ "format": "ND",
332+ "shape": [
333+ -2
334+ ],
335+ "paramType": "required"
336+ },
337+ {
338+ "name": "key",
339+ "index": 1,
340+ "dtype": "uint64",
341+ "format": "ND",
342+ "shape": [
343+ -2
344+ ],
345+ "paramType": "required"
346+ },
347+ {
348+ "name": "counter",
349+ "index": 2,
350+ "dtype": "uint64",
351+ "format": "ND",
352+ "shape": [
353+ -2
354+ ],
355+ "paramType": "required"
356+ },
357+ {
358+ "name": "alg",
359+ "index": 3,
360+ "dtype": "int32",
361+ "format": "ND",
362+ "shape": [
363+ -2
364+ ],
365+ "paramType": "required"
366+ }
367+ ],
368+ "outputs": [
369+ {
370+ "name": "y",
371+ "index": 0,
372+ "dtype": "bfloat16",
373+ "format": "ND",
374+ "shape": [
375+ -2
376+ ],
377+ "paramType": "required"
378+ }
379+ ],
380+ "attrs": [
381+ {
382+ "name": "dtype",
383+ "dtype": "int",
384+ "value": null
385+ }
386+ ]
387+ }
388+ ]
389+}
Arandom/stateless_random_normal_v2/op_host/config/ascend960pr/stateless_random_normal_v2_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，ascnedxx为可选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+[StatelessRandomNormalV2]
13+default=0
Mrandom/stateless_random_normal_v2/op_host/stateless_random_normal_v2_def.cpp+2-0
@@ -56,6 +56,8 @@ public:
56 this->Attr("dtype").AttrType(OPTIONAL).Int(0);56 this->Attr("dtype").AttrType(OPTIONAL).Int(0);
57 this->AICore().AddConfig("ascend950");57 this->AICore().AddConfig("ascend950");
58 this->AICore().AddConfig("ascend350");58 this->AICore().AddConfig("ascend350");
59+ this->AICore().AddConfig("ascend960dt");
60+ this->AICore().AddConfig("ascend960pr");
59 }61 }
60};62};
61 63 
Mrandom/stateless_random_normal_v3/CMakeLists.txt+2-2
@@ -9,7 +9,7 @@
9# ---------------------------------------------------------------------------------------------------------9# ---------------------------------------------------------------------------------------------------------
10 10 
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12-set(SUPPORT_COMPUTE_UNIT "ascend950" "ascend350")12+set(SUPPORT_COMPUTE_UNIT "ascend950" "ascend350" "ascend960dt" "ascend960pr")
13# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch35" "arch35" "arch35" "arch35")
15add_all_modules_sources(OPTYPE stateless_random_normal_v3 ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DEPENDENCIES random_common DISABLE_IN_OPP TRUE)15add_all_modules_sources(OPTYPE stateless_random_normal_v3 ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DEPENDENCIES random_common DISABLE_IN_OPP TRUE)
Mrandom/stateless_random_normal_v3/op_host/arch35/stateless_random_normal_v3_tiling_arch35.cpp+47-24
@@ -9,14 +9,14 @@
9 */9 */
10 10 
11/*!11/*!
12-* \file stateless_random_normal_v3_tiling_arch35.cpp12+ * \file stateless_random_normal_v3_tiling_arch35.cpp
13-* \brief13+ * \brief
14-*/14+ */
15 15 
16#include "platform/platform_infos_def.h"16#include "platform/platform_infos_def.h"
17#include "platform/platform_ascendc.h"17#include "platform/platform_ascendc.h"
18#include "op_common/op_host/util/platform_util.h"18#include "op_common/op_host/util/platform_util.h"
19-#include "../../../random_common/op_host/arch35/random_tiling_base.h"19+#include "../../../random_common/op_host/arch35/random_tiling_base.h"
20#include "exe_graph/runtime/shape.h"20#include "exe_graph/runtime/shape.h"
21#include "op_host/tiling_base_class.h"21#include "op_host/tiling_base_class.h"
22#include "stateless_random_normal_v3_tiling_arch35.h"22#include "stateless_random_normal_v3_tiling_arch35.h"
@@ -32,6 +32,7 @@ static constexpr uint16_t INPUT_IDX_MEAN = 3;
32static constexpr uint16_t INPUT_IDX_STDEV = 4;32static constexpr uint16_t INPUT_IDX_STDEV = 4;
33static constexpr uint16_t OUTPUT_IDX_Y = 0;33static constexpr uint16_t OUTPUT_IDX_Y = 0;
34static constexpr uint16_t CORE_ALIGN_SIZE = 512;34static constexpr uint16_t CORE_ALIGN_SIZE = 512;
35+static constexpr uint16_t CORE_ALIGN_SIZE_ARCH92 = 32;
35static constexpr uint32_t MEAN_SCALAR_FLAG = 1;36static constexpr uint32_t MEAN_SCALAR_FLAG = 1;
36static constexpr uint32_t STDEV_SCALAR_FLAG = 2;37static constexpr uint32_t STDEV_SCALAR_FLAG = 2;
37static constexpr uint32_t BASE_SPLIT_UB_NUM = 5;38static constexpr uint32_t BASE_SPLIT_UB_NUM = 5;
@@ -40,16 +41,15 @@ OpTilingConfig StatelessRandomNormalV3Tiling::BuildOpConfig()
40{41{
41 OpTilingConfig config;42 OpTilingConfig config;
42 config.inputCheckRules = {43 config.inputCheckRules = {
43- // 输入索引: dtype列表，shapeSize，dim_num44+ // 输入索引: dtype列表，shapeSize，dim_num
44- {0, {{ge::DT_INT32, ge::DT_INT64}, -1, {1}, nullptr}}, // shape45+ {0, {{ge::DT_INT32, ge::DT_INT64}, -1, {1}, nullptr}}, // shape
45- {1, {{ge::DT_UINT64}, 1, {}, nullptr}}, // key46+ {1, {{ge::DT_UINT64}, 1, {}, nullptr}}, // key
46- {2, {{ge::DT_UINT64}, -1, {}, nullptr}}, // counter47+ {2, {{ge::DT_UINT64}, -1, {}, nullptr}}, // counter
47- {3, {{ge::DT_FLOAT}, -1, {}, nullptr}}, // mean (scalar or tensor)48+ {3, {{ge::DT_FLOAT}, -1, {}, nullptr}}, // mean (scalar or tensor)
48- {4, {{ge::DT_FLOAT}, -1, {}, nullptr}}, // stdev (scalar or tensor)49+ {4, {{ge::DT_FLOAT}, -1, {}, nullptr}}, // stdev (scalar or tensor)
49- };50+ };
50 config.outputCheckRules = {51 config.outputCheckRules = {
51- {0, {{ge::DT_FLOAT, ge::DT_BF16, ge::DT_FLOAT16}, -1, {0,1,2,3,4,5,6,7,8}, nullptr}}52+ {0, {{ge::DT_FLOAT, ge::DT_BF16, ge::DT_FLOAT16}, -1, {0, 1, 2, 3, 4, 5, 6, 7, 8}, nullptr}}}; // y
52- }; // y
53 config.getOutputSize = [](gert::TilingContext* ctx, int64_t& shapeSize) -> ge::graphStatus {53 config.getOutputSize = [](gert::TilingContext* ctx, int64_t& shapeSize) -> ge::graphStatus {
54 auto outputShape = ctx->GetOutputShape(OUTPUT_IDX_Y);54 auto outputShape = ctx->GetOutputShape(OUTPUT_IDX_Y);
55 OP_CHECK_NULL_WITH_CONTEXT(ctx, outputShape);55 OP_CHECK_NULL_WITH_CONTEXT(ctx, outputShape);
@@ -57,9 +57,14 @@ OpTilingConfig StatelessRandomNormalV3Tiling::BuildOpConfig()
57 return ge::GRAPH_SUCCESS;57 return ge::GRAPH_SUCCESS;
58 };58 };
59 59 
60- config.getKeyAndCounter = []([[maybe_unused]] gert::TilingContext* ctx, uint32_t key[2], uint32_t counter[4]) -> ge::graphStatus {60+ config.getKeyAndCounter = []([[maybe_unused]] gert::TilingContext* ctx, uint32_t key[2],
61- key[0] = 0; key[1] = 0;61+ uint32_t counter[4]) -> ge::graphStatus {
62- counter[0] = 0; counter[1] = 0; counter[2] = 0; counter[3] = 0;62+ key[0] = 0;
63+ key[1] = 0;
64+ counter[0] = 0;
65+ counter[1] = 0;
66+ counter[2] = 0;
67+ counter[3] = 0;
63 return ge::GRAPH_SUCCESS;68 return ge::GRAPH_SUCCESS;
64 };69 };
65 70 
@@ -70,12 +75,12 @@ OpTilingConfig StatelessRandomNormalV3Tiling::BuildOpConfig()
70 auto meanTensor = ctx->GetInputTensor(INPUT_IDX_MEAN);75 auto meanTensor = ctx->GetInputTensor(INPUT_IDX_MEAN);
71 OP_CHECK_NULL_WITH_CONTEXT(ctx, meanTensor);76 OP_CHECK_NULL_WITH_CONTEXT(ctx, meanTensor);
72 if (meanTensor->GetShapeSize() > 1) {77 if (meanTensor->GetShapeSize() > 1) {
73- extraBufNum++; 78+ extraBufNum++;
74 }79 }
75 auto stdevTensor = ctx->GetInputTensor(INPUT_IDX_STDEV);80 auto stdevTensor = ctx->GetInputTensor(INPUT_IDX_STDEV);
76 OP_CHECK_NULL_WITH_CONTEXT(ctx, stdevTensor);81 OP_CHECK_NULL_WITH_CONTEXT(ctx, stdevTensor);
77 if (stdevTensor->GetShapeSize() > 1) {82 if (stdevTensor->GetShapeSize() > 1) {
78- extraBufNum++; 83+ extraBufNum++;
79 }84 }
80 85 
81 bufNum = sizeof(float) * (BASE_SPLIT_UB_NUM + extraBufNum);86 bufNum = sizeof(float) * (BASE_SPLIT_UB_NUM + extraBufNum);
@@ -87,7 +92,23 @@ OpTilingConfig StatelessRandomNormalV3Tiling::BuildOpConfig()
87 return config;92 return config;
88}93}
89 94 
90-StatelessRandomNormalV3Tiling::StatelessRandomNormalV3Tiling(gert::TilingContext* ctx) : RandomTilingArch35(ctx, BuildOpConfig()){}95+StatelessRandomNormalV3Tiling::StatelessRandomNormalV3Tiling(gert::TilingContext* ctx)
96+ : RandomTilingArch35(ctx, BuildOpConfig())
97+{}
98+ 
99+ge::graphStatus StatelessRandomNormalV3Tiling::BeforeProcess()
100+{
101+ auto platformInfo = context_->GetPlatformInfo();
102+ if (platformInfo == nullptr) {
103+ return ge::GRAPH_SUCCESS;
104+ }
105+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
106+ auto npuArch = ascendcPlatform.GetCurNpuArch();
107+ if (npuArch == NpuArch::DAV_9201 || npuArch == NpuArch::DAV_9202) {
108+ config_.coreAlignSize = CORE_ALIGN_SIZE_ARCH92;
109+ }
110+ return ge::GRAPH_SUCCESS;
111+}
91 112 
92ge::graphStatus StatelessRandomNormalV3Tiling::UniqueProcess()113ge::graphStatus StatelessRandomNormalV3Tiling::UniqueProcess()
93{114{
@@ -104,9 +125,10 @@ ge::graphStatus StatelessRandomNormalV3Tiling::UniqueProcess()
104 v3KernelMode |= MEAN_SCALAR_FLAG;125 v3KernelMode |= MEAN_SCALAR_FLAG;
105 } else if (meanSize != outputSize) {126 } else if (meanSize != outputSize) {
106 std::string valueStr = std::to_string(meanSize) + " and " + std::to_string(outputSize);127 std::string valueStr = std::to_string(meanSize) + " and " + std::to_string(outputSize);
107- std::string reasonMsg = "StatelessRandomNormalV3 does not support broadcast for mean, mean shapeSize must equal output shapeSize or be 1";128+ std::string reasonMsg = "StatelessRandomNormalV3 does not support broadcast for mean, mean shapeSize must "
129+ "equal output shapeSize or be 1";
108 OP_LOGE_FOR_INVALID_SHAPESIZES_WITH_REASON(context_->GetNodeName(), "mean shapeSize and output shapeSize",130 OP_LOGE_FOR_INVALID_SHAPESIZES_WITH_REASON(context_->GetNodeName(), "mean shapeSize and output shapeSize",
109- valueStr.c_str(), reasonMsg.c_str());131+ valueStr.c_str(), reasonMsg.c_str());
110 return ge::GRAPH_FAILED;132 return ge::GRAPH_FAILED;
111 }133 }
112 134 
@@ -117,9 +139,10 @@ ge::graphStatus StatelessRandomNormalV3Tiling::UniqueProcess()
117 v3KernelMode |= STDEV_SCALAR_FLAG;139 v3KernelMode |= STDEV_SCALAR_FLAG;
118 } else if (stdevSize != outputSize) {140 } else if (stdevSize != outputSize) {
119 std::string valueStr = std::to_string(stdevSize) + " and " + std::to_string(outputSize);141 std::string valueStr = std::to_string(stdevSize) + " and " + std::to_string(outputSize);
120- std::string reasonMsg = "StatelessRandomNormalV3 does not support broadcast for stdev, stdev shapeSize must equal output shapeSize or be 1";142+ std::string reasonMsg = "StatelessRandomNormalV3 does not support broadcast for stdev, stdev shapeSize must "
143+ "equal output shapeSize or be 1";
121 OP_LOGE_FOR_INVALID_SHAPESIZES_WITH_REASON(context_->GetNodeName(), "stdev shapeSize and output shapeSize",144 OP_LOGE_FOR_INVALID_SHAPESIZES_WITH_REASON(context_->GetNodeName(), "stdev shapeSize and output shapeSize",
122- valueStr.c_str(), reasonMsg.c_str());145+ valueStr.c_str(), reasonMsg.c_str());
123 return ge::GRAPH_FAILED;146 return ge::GRAPH_FAILED;
124 }147 }
125 148 
@@ -143,4 +166,4 @@ IMPL_OP_OPTILING(StatelessRandomNormalV3)
143 .Tiling(TilingStatelessRandomNormalV3)166 .Tiling(TilingStatelessRandomNormalV3)
144 .TilingParse<RandomOperatorCompileInfo>(TilingPrepare4StatelessRandomNormalV3Tiling);167 .TilingParse<RandomOperatorCompileInfo>(TilingPrepare4StatelessRandomNormalV3Tiling);
145 168 
146-} // namespace optiling169+} // namespace optiling
Mrandom/stateless_random_normal_v3/op_host/arch35/stateless_random_normal_v3_tiling_arch35.h+7-7
@@ -9,9 +9,9 @@
9 */9 */
10 10 
11/*!11/*!
12-* \file stateless_random_normal_v3_tiling.h12+ * \file stateless_random_normal_v3_tiling.h
13-* \brief13+ * \brief
14-*/14+ */
15 15 
16#ifndef STATELESS_RANDOM_NORMAL_V3_TILING_ARCH35_H16#ifndef STATELESS_RANDOM_NORMAL_V3_TILING_ARCH35_H
17#define STATELESS_RANDOM_NORMAL_V3_TILING_ARCH35_H17#define STATELESS_RANDOM_NORMAL_V3_TILING_ARCH35_H
@@ -23,16 +23,16 @@
23#include "tiling/platform/platform_ascendc.h"23#include "tiling/platform/platform_ascendc.h"
24#include "log/log.h"24#include "log/log.h"
25#include "util/math_util.h"25#include "util/math_util.h"
26-#include "../../../random_common/op_host/arch35/random_tiling_arch35.h"26+#include "../../../random_common/op_host/arch35/random_tiling_arch35.h"
27 27 
28namespace optiling {28namespace optiling {
29-class StatelessRandomNormalV3Tiling : public RandomTilingArch3529+class StatelessRandomNormalV3Tiling : public RandomTilingArch35 {
30-{
31public:30public:
32 explicit StatelessRandomNormalV3Tiling(gert::TilingContext* ctx);31 explicit StatelessRandomNormalV3Tiling(gert::TilingContext* ctx);
33 ~StatelessRandomNormalV3Tiling() override = default;32 ~StatelessRandomNormalV3Tiling() override = default;
34 33 
35protected:34protected:
35+ ge::graphStatus BeforeProcess() override;
36 ge::graphStatus UniqueProcess() override;36 ge::graphStatus UniqueProcess() override;
37 37 
38private:38private:
@@ -40,4 +40,4 @@ private:
40};40};
41 41 
42} // namespace optiling42} // namespace optiling
43-#endif // STATELESS_RANDOM_NORMAL_V3_TILING_ARCH35_H43+#endif // STATELESS_RANDOM_NORMAL_V3_TILING_ARCH35_H
Arandom/stateless_random_normal_v3/op_host/config/ascend960dt/stateless_random_normal_v3_binary.json+449-0
@@ -0,0 +1,449 @@
1+{
2+ "op_type": "StatelessRandomNormalV3",
3+ "op_list": [
4+ {
5+ "bin_filename": "StatelessRandomNormalV3_1ea4bdp73cd9df67072d9z7aaec9c583",
6+ "inputs": [
7+ {
8+ "name": "shape",
9+ "index": 0,
10+ "dtype": "int32",
11+ "format": "ND",
12+ "shape": [
13+ -2
14+ ],
15+ "paramType": "required"
16+ },
17+ {
18+ "name": "key",
19+ "index": 1,
20+ "dtype": "uint64",
21+ "format": "ND",
22+ "shape": [
23+ -2
24+ ],
25+ "paramType": "required"
26+ },
27+ {
28+ "name": "counter",
29+ "index": 2,
30+ "dtype": "uint64",
31+ "format": "ND",
32+ "shape": [
33+ -2
34+ ],
35+ "paramType": "required"
36+ },
37+ {
38+ "name": "mean",
39+ "index": 3,
40+ "dtype": "float32",
41+ "format": "ND",
42+ "shape": [
43+ -2
44+ ],
45+ "paramType": "required"
46+ },
47+ {
48+ "name": "stdev",
49+ "index": 4,
50+ "dtype": "float32",
51+ "format": "ND",
52+ "shape": [
53+ -2
54+ ],
55+ "paramType": "required"
56+ }
57+ ],
58+ "outputs": [
59+ {
60+ "name": "y",
61+ "index": 0,
62+ "dtype": "float32",
63+ "format": "ND",
64+ "shape": [
65+ -2
66+ ],
67+ "paramType": "required"
68+ }
69+ ],
70+ "attrs": [
71+ {
72+ "name": "dtype",
73+ "dtype": "int",
74+ "value": null
75+ }
76+ ]
77+ },
78+ {
79+ "bin_filename": "StatelessRandomNormalV3_1ea4bd0e3cd9df67072d9z7akyc9c583",
80+ "inputs": [
81+ {
82+ "name": "shape",
83+ "index": 0,
84+ "dtype": "int32",
85+ "format": "ND",
86+ "shape": [
87+ -2
88+ ],
89+ "paramType": "required"
90+ },
91+ {
92+ "name": "key",
93+ "index": 1,
94+ "dtype": "uint64",
95+ "format": "ND",
96+ "shape": [
97+ -2
98+ ],
99+ "paramType": "required"
100+ },
101+ {
102+ "name": "counter",
103+ "index": 2,
104+ "dtype": "uint64",
105+ "format": "ND",
106+ "shape": [
107+ -2
108+ ],
109+ "paramType": "required"
110+ },
111+ {
112+ "name": "mean",
113+ "index": 3,
114+ "dtype": "float32",
115+ "format": "ND",
116+ "shape": [
117+ -2
118+ ],
119+ "paramType": "required"
120+ },
121+ {
122+ "name": "stdev",
123+ "index": 4,
124+ "dtype": "float32",
125+ "format": "ND",
126+ "shape": [
127+ -2
128+ ],
129+ "paramType": "required"
130+ }
131+ ],
132+ "outputs": [
133+ {
134+ "name": "y",
135+ "index": 0,
136+ "dtype": "float16",
137+ "format": "ND",
138+ "shape": [
139+ -2
140+ ],
141+ "paramType": "required"
142+ }
143+ ],
144+ "attrs": [
145+ {
146+ "name": "dtype",
147+ "dtype": "int",
148+ "value": null
149+ }
150+ ]
151+ },
152+ {
153+ "bin_filename": "StatelessRandomNormalV3_1ea4bd9g3cd9df67072d9z7amgc9c733",
154+ "inputs": [
155+ {
156+ "name": "shape",
157+ "index": 0,
158+ "dtype": "int32",
159+ "format": "ND",
160+ "shape": [
161+ -2
162+ ],
163+ "paramType": "required"
164+ },
165+ {
166+ "name": "key",
167+ "index": 1,
168+ "dtype": "uint64",
169+ "format": "ND",
170+ "shape": [
171+ -2
172+ ],
173+ "paramType": "required"
174+ },
175+ {
176+ "name": "counter",
177+ "index": 2,
178+ "dtype": "uint64",
179+ "format": "ND",
180+ "shape": [
181+ -2
182+ ],
183+ "paramType": "required"
184+ },
185+ {
186+ "name": "mean",
187+ "index": 3,
188+ "dtype": "float32",
189+ "format": "ND",
190+ "shape": [
191+ -2
192+ ],
193+ "paramType": "required"
194+ },
195+ {
196+ "name": "stdev",
197+ "index": 4,
198+ "dtype": "float32",
199+ "format": "ND",
200+ "shape": [
201+ -2
202+ ],
203+ "paramType": "required"
204+ }
205+ ],
206+ "outputs": [
207+ {
208+ "name": "y",
209+ "index": 0,
210+ "dtype": "bfloat16",
211+ "format": "ND",
212+ "shape": [
213+ -2
214+ ],
215+ "paramType": "required"
216+ }
217+ ],
218+ "attrs": [
219+ {
220+ "name": "dtype",
221+ "dtype": "int",
222+ "value": null
223+ }
224+ ]
225+ },
226+ {
227+ "bin_filename": "StatelessRandomNormalV3_1ea4bd0e3cd9df67072d9z7aaec9c583",
228+ "inputs": [
229+ {
230+ "name": "shape",
231+ "index": 0,
232+ "dtype": "int64",
233+ "format": "ND",
234+ "shape": [
235+ -2
236+ ],
237+ "paramType": "required"
238+ },
239+ {
240+ "name": "key",
241+ "index": 1,
242+ "dtype": "uint64",
243+ "format": "ND",
244+ "shape": [
245+ -2
246+ ],
247+ "paramType": "required"
248+ },
249+ {
250+ "name": "counter",
251+ "index": 2,
252+ "dtype": "uint64",
253+ "format": "ND",
254+ "shape": [
255+ -2
256+ ],
257+ "paramType": "required"
258+ },
259+ {
260+ "name": "mean",
261+ "index": 3,
262+ "dtype": "float32",
263+ "format": "ND",
264+ "shape": [
265+ -2
266+ ],
267+ "paramType": "required"
268+ },
269+ {
270+ "name": "stdev",
271+ "index": 4,
272+ "dtype": "float32",
273+ "format": "ND",
274+ "shape": [
275+ -2
276+ ],
277+ "paramType": "required"
278+ }
279+ ],
280+ "outputs": [
281+ {
282+ "name": "y",
283+ "index": 0,
284+ "dtype": "float32",
285+ "format": "ND",
286+ "shape": [
287+ -2
288+ ],
289+ "paramType": "required"
290+ }
291+ ],
292+ "attrs": [
293+ {
294+ "name": "dtype",
295+ "dtype": "int",
296+ "value": null
297+ }
298+ ]
299+ },
300+ {
301+ "bin_filename": "StatelessRandomNormalV3_1ea4bd0e3cd9df67072d9z7amgc9c583",
302+ "inputs": [
303+ {
304+ "name": "shape",
305+ "index": 0,
306+ "dtype": "int64",
307+ "format": "ND",
308+ "shape": [
309+ -2
310+ ],
311+ "paramType": "required"
312+ },
313+ {
314+ "name": "key",
315+ "index": 1,
316+ "dtype": "uint64",
317+ "format": "ND",
318+ "shape": [
319+ -2
320+ ],
321+ "paramType": "required"
322+ },
323+ {
324+ "name": "counter",
325+ "index": 2,
326+ "dtype": "uint64",
327+ "format": "ND",
328+ "shape": [
329+ -2
330+ ],
331+ "paramType": "required"
332+ },
333+ {
334+ "name": "mean",
335+ "index": 3,
336+ "dtype": "float32",
337+ "format": "ND",
338+ "shape": [
339+ -2
340+ ],
341+ "paramType": "required"
342+ },
343+ {
344+ "name": "stdev",
345+ "index": 4,
346+ "dtype": "float32",
347+ "format": "ND",
348+ "shape": [
349+ -2
350+ ],
351+ "paramType": "required"
352+ }
353+ ],
354+ "outputs": [
355+ {
356+ "name": "y",
357+ "index": 0,
358+ "dtype": "float16",
359+ "format": "ND",
360+ "shape": [
361+ -2
362+ ],
363+ "paramType": "required"
364+ }
365+ ],
366+ "attrs": [
367+ {
368+ "name": "dtype",
369+ "dtype": "int",
370+ "value": null
371+ }
372+ ]
373+ },
374+ {
375+ "bin_filename": "StatelessRandomNormalV3_1ea4bd0e3cd9df67072d9z7amgc9c733",
376+ "inputs": [
377+ {
378+ "name": "shape",
379+ "index": 0,
380+ "dtype": "int64",
381+ "format": "ND",
382+ "shape": [
383+ -2
384+ ],
385+ "paramType": "required"
386+ },
387+ {
388+ "name": "key",
389+ "index": 1,
390+ "dtype": "uint64",
391+ "format": "ND",
392+ "shape": [
393+ -2
394+ ],
395+ "paramType": "required"
396+ },
397+ {
398+ "name": "counter",
399+ "index": 2,
400+ "dtype": "uint64",
401+ "format": "ND",
402+ "shape": [
403+ -2
404+ ],
405+ "paramType": "required"
406+ },
407+ {
408+ "name": "mean",
409+ "index": 3,
410+ "dtype": "float32",
411+ "format": "ND",
412+ "shape": [
413+ -2
414+ ],
415+ "paramType": "required"
416+ },
417+ {
418+ "name": "stdev",
419+ "index": 4,
420+ "dtype": "float32",
421+ "format": "ND",
422+ "shape": [
423+ -2
424+ ],
425+ "paramType": "required"
426+ }
427+ ],
428+ "outputs": [
429+ {
430+ "name": "y",
431+ "index": 0,
432+ "dtype": "bfloat16",
433+ "format": "ND",
434+ "shape": [
435+ -2
436+ ],
437+ "paramType": "required"
438+ }
439+ ],
440+ "attrs": [
441+ {
442+ "name": "dtype",
443+ "dtype": "int",
444+ "value": null
445+ }
446+ ]
447+ }
448+ ]
449+}
Arandom/stateless_random_normal_v3/op_host/config/ascend960dt/stateless_random_normal_v3_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，ascnedxx为可选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+[StatelessRandomNormalV3]
13+default=0
Arandom/stateless_random_normal_v3/op_host/config/ascend960pr/stateless_random_normal_v3_binary.json+449-0
@@ -0,0 +1,449 @@
1+{
2+ "op_type": "StatelessRandomNormalV3",
3+ "op_list": [
4+ {
5+ "bin_filename": "StatelessRandomNormalV3_1ea4bdp73cd9df67072d9z7aaec9c583",
6+ "inputs": [
7+ {
8+ "name": "shape",
9+ "index": 0,
10+ "dtype": "int32",
11+ "format": "ND",
12+ "shape": [
13+ -2
14+ ],
15+ "paramType": "required"
16+ },
17+ {
18+ "name": "key",
19+ "index": 1,
20+ "dtype": "uint64",
21+ "format": "ND",
22+ "shape": [
23+ -2
24+ ],
25+ "paramType": "required"
26+ },
27+ {
28+ "name": "counter",
29+ "index": 2,
30+ "dtype": "uint64",
31+ "format": "ND",
32+ "shape": [
33+ -2
34+ ],
35+ "paramType": "required"
36+ },
37+ {
38+ "name": "mean",
39+ "index": 3,
40+ "dtype": "float32",
41+ "format": "ND",
42+ "shape": [
43+ -2
44+ ],
45+ "paramType": "required"
46+ },
47+ {
48+ "name": "stdev",
49+ "index": 4,
50+ "dtype": "float32",
51+ "format": "ND",
52+ "shape": [
53+ -2
54+ ],
55+ "paramType": "required"
56+ }
57+ ],
58+ "outputs": [
59+ {
60+ "name": "y",
61+ "index": 0,
62+ "dtype": "float32",
63+ "format": "ND",
64+ "shape": [
65+ -2
66+ ],
67+ "paramType": "required"
68+ }
69+ ],
70+ "attrs": [
71+ {
72+ "name": "dtype",
73+ "dtype": "int",
74+ "value": null
75+ }
76+ ]
77+ },
78+ {
79+ "bin_filename": "StatelessRandomNormalV3_1ea4bd0e3cd9df67072d9z7akyc9c583",
80+ "inputs": [
81+ {
82+ "name": "shape",
83+ "index": 0,
84+ "dtype": "int32",
85+ "format": "ND",
86+ "shape": [
87+ -2
88+ ],
89+ "paramType": "required"
90+ },
91+ {
92+ "name": "key",
93+ "index": 1,
94+ "dtype": "uint64",
95+ "format": "ND",
96+ "shape": [
97+ -2
98+ ],
99+ "paramType": "required"
100+ },
101+ {
102+ "name": "counter",
103+ "index": 2,
104+ "dtype": "uint64",
105+ "format": "ND",
106+ "shape": [
107+ -2
108+ ],
109+ "paramType": "required"
110+ },
111+ {
112+ "name": "mean",
113+ "index": 3,
114+ "dtype": "float32",
115+ "format": "ND",
116+ "shape": [
117+ -2
118+ ],
119+ "paramType": "required"
120+ },
121+ {
122+ "name": "stdev",
123+ "index": 4,
124+ "dtype": "float32",
125+ "format": "ND",
126+ "shape": [
127+ -2
128+ ],
129+ "paramType": "required"
130+ }
131+ ],
132+ "outputs": [
133+ {
134+ "name": "y",
135+ "index": 0,
136+ "dtype": "float16",
137+ "format": "ND",
138+ "shape": [
139+ -2
140+ ],
141+ "paramType": "required"
142+ }
143+ ],
144+ "attrs": [
145+ {
146+ "name": "dtype",
147+ "dtype": "int",
148+ "value": null
149+ }
150+ ]
151+ },
152+ {
153+ "bin_filename": "StatelessRandomNormalV3_1ea4bd9g3cd9df67072d9z7amgc9c733",
154+ "inputs": [
155+ {
156+ "name": "shape",
157+ "index": 0,
158+ "dtype": "int32",
159+ "format": "ND",
160+ "shape": [
161+ -2
162+ ],
163+ "paramType": "required"
164+ },
165+ {
166+ "name": "key",
167+ "index": 1,
168+ "dtype": "uint64",
169+ "format": "ND",
170+ "shape": [
171+ -2
172+ ],
173+ "paramType": "required"
174+ },
175+ {
176+ "name": "counter",
177+ "index": 2,
178+ "dtype": "uint64",
179+ "format": "ND",
180+ "shape": [
181+ -2
182+ ],
183+ "paramType": "required"
184+ },
185+ {
186+ "name": "mean",
187+ "index": 3,
188+ "dtype": "float32",
189+ "format": "ND",
190+ "shape": [
191+ -2
192+ ],
193+ "paramType": "required"
194+ },
195+ {
196+ "name": "stdev",
197+ "index": 4,
198+ "dtype": "float32",
199+ "format": "ND",
200+ "shape": [
201+ -2
202+ ],
203+ "paramType": "required"
204+ }
205+ ],
206+ "outputs": [
207+ {
208+ "name": "y",
209+ "index": 0,
210+ "dtype": "bfloat16",
211+ "format": "ND",
212+ "shape": [
213+ -2
214+ ],
215+ "paramType": "required"
216+ }
217+ ],
218+ "attrs": [
219+ {
220+ "name": "dtype",
221+ "dtype": "int",
222+ "value": null
223+ }
224+ ]
225+ },
226+ {
227+ "bin_filename": "StatelessRandomNormalV3_1ea4bd0e3cd9df67072d9z7aaec9c583",
228+ "inputs": [
229+ {
230+ "name": "shape",
231+ "index": 0,
232+ "dtype": "int64",
233+ "format": "ND",
234+ "shape": [
235+ -2
236+ ],
237+ "paramType": "required"
238+ },
239+ {
240+ "name": "key",
241+ "index": 1,
242+ "dtype": "uint64",
243+ "format": "ND",
244+ "shape": [
245+ -2
246+ ],
247+ "paramType": "required"
248+ },
249+ {
250+ "name": "counter",
251+ "index": 2,
252+ "dtype": "uint64",
253+ "format": "ND",
254+ "shape": [
255+ -2
256+ ],
257+ "paramType": "required"
258+ },
259+ {
260+ "name": "mean",
261+ "index": 3,
262+ "dtype": "float32",
263+ "format": "ND",
264+ "shape": [
265+ -2
266+ ],
267+ "paramType": "required"
268+ },
269+ {
270+ "name": "stdev",
271+ "index": 4,
272+ "dtype": "float32",
273+ "format": "ND",
274+ "shape": [
275+ -2
276+ ],
277+ "paramType": "required"
278+ }
279+ ],
280+ "outputs": [
281+ {
282+ "name": "y",
283+ "index": 0,
284+ "dtype": "float32",
285+ "format": "ND",
286+ "shape": [
287+ -2
288+ ],
289+ "paramType": "required"
290+ }
291+ ],
292+ "attrs": [
293+ {
294+ "name": "dtype",
295+ "dtype": "int",
296+ "value": null
297+ }
298+ ]
299+ },
300+ {
301+ "bin_filename": "StatelessRandomNormalV3_1ea4bd0e3cd9df67072d9z7amgc9c583",
302+ "inputs": [
303+ {
304+ "name": "shape",
305+ "index": 0,
306+ "dtype": "int64",
307+ "format": "ND",
308+ "shape": [
309+ -2
310+ ],
311+ "paramType": "required"
312+ },
313+ {
314+ "name": "key",
315+ "index": 1,
316+ "dtype": "uint64",
317+ "format": "ND",
318+ "shape": [
319+ -2
320+ ],
321+ "paramType": "required"
322+ },
323+ {
324+ "name": "counter",
325+ "index": 2,
326+ "dtype": "uint64",
327+ "format": "ND",
328+ "shape": [
329+ -2
330+ ],
331+ "paramType": "required"
332+ },
333+ {
334+ "name": "mean",
335+ "index": 3,
336+ "dtype": "float32",
337+ "format": "ND",
338+ "shape": [
339+ -2
340+ ],
341+ "paramType": "required"
342+ },
343+ {
344+ "name": "stdev",
345+ "index": 4,
346+ "dtype": "float32",
347+ "format": "ND",
348+ "shape": [
349+ -2
350+ ],
351+ "paramType": "required"
352+ }
353+ ],
354+ "outputs": [
355+ {
356+ "name": "y",
357+ "index": 0,
358+ "dtype": "float16",
359+ "format": "ND",
360+ "shape": [
361+ -2
362+ ],
363+ "paramType": "required"
364+ }
365+ ],
366+ "attrs": [
367+ {
368+ "name": "dtype",
369+ "dtype": "int",
370+ "value": null
371+ }
372+ ]
373+ },
374+ {
375+ "bin_filename": "StatelessRandomNormalV3_1ea4bd0e3cd9df67072d9z7amgc9c733",
376+ "inputs": [
377+ {
378+ "name": "shape",
379+ "index": 0,
380+ "dtype": "int64",
381+ "format": "ND",
382+ "shape": [
383+ -2
384+ ],
385+ "paramType": "required"
386+ },
387+ {
388+ "name": "key",
389+ "index": 1,
390+ "dtype": "uint64",
391+ "format": "ND",
392+ "shape": [
393+ -2
394+ ],
395+ "paramType": "required"
396+ },
397+ {
398+ "name": "counter",
399+ "index": 2,
400+ "dtype": "uint64",
401+ "format": "ND",
402+ "shape": [
403+ -2
404+ ],
405+ "paramType": "required"
406+ },
407+ {
408+ "name": "mean",
409+ "index": 3,
410+ "dtype": "float32",
411+ "format": "ND",
412+ "shape": [
413+ -2
414+ ],
415+ "paramType": "required"
416+ },
417+ {
418+ "name": "stdev",
419+ "index": 4,
420+ "dtype": "float32",
421+ "format": "ND",
422+ "shape": [
423+ -2
424+ ],
425+ "paramType": "required"
426+ }
427+ ],
428+ "outputs": [
429+ {
430+ "name": "y",
431+ "index": 0,
432+ "dtype": "bfloat16",
433+ "format": "ND",
434+ "shape": [
435+ -2
436+ ],
437+ "paramType": "required"
438+ }
439+ ],
440+ "attrs": [
441+ {
442+ "name": "dtype",
443+ "dtype": "int",
444+ "value": null
445+ }
446+ ]
447+ }
448+ ]
449+}
Arandom/stateless_random_normal_v3/op_host/config/ascend960pr/stateless_random_normal_v3_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，ascnedxx为可选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+[StatelessRandomNormalV3]
13+default=0
Mrandom/stateless_random_normal_v3/op_host/stateless_random_normal_v3_def.cpp+2-0
@@ -63,6 +63,8 @@ public:
63 this->Attr("dtype").AttrType(OPTIONAL).Int(0);63 this->Attr("dtype").AttrType(OPTIONAL).Int(0);
64 this->AICore().AddConfig("ascend950");64 this->AICore().AddConfig("ascend950");
65 this->AICore().AddConfig("ascend350");65 this->AICore().AddConfig("ascend350");
66+ this->AICore().AddConfig("ascend960dt");
67+ this->AICore().AddConfig("ascend960pr");
66 }68 }
67};69};
68 70 
Mrandom/stateless_random_normal_v3/op_kernel/stateless_random_normal_v3.cpp+6-4
@@ -16,9 +16,10 @@ namespace {
16#define STATELESSRANDOMNORMALV3_TILING_KEY 10016#define STATELESSRANDOMNORMALV3_TILING_KEY 100
17} // namespace17} // namespace
18 18 
19-extern "C" __global__ __aicore__ void stateless_random_normal_v3(19+#if defined(__DAV_VEC__)
20- GM_ADDR shape, GM_ADDR key, GM_ADDR counter, GM_ADDR mean, GM_ADDR stdev, GM_ADDR y, GM_ADDR workspace,20+extern "C" __global__ __aicore__ void stateless_random_normal_v3(GM_ADDR shape, GM_ADDR key, GM_ADDR counter,
21- GM_ADDR tiling)21+ GM_ADDR mean, GM_ADDR stdev, GM_ADDR y,
22+ GM_ADDR workspace, GM_ADDR tiling)
22{23{
23 REGISTER_TILING_DEFAULT(RandomUnifiedTilingDataStruct);24 REGISTER_TILING_DEFAULT(RandomUnifiedTilingDataStruct);
24 GET_TILING_DATA(tilingData, tiling);25 GET_TILING_DATA(tilingData, tiling);
@@ -30,4 +31,5 @@ extern "C" __global__ __aicore__ void stateless_random_normal_v3(
30 op.Init(y, key, counter, mean, stdev);31 op.Init(y, key, counter, mean, stdev);
31 op.Process();32 op.Process();
32 }33 }
33-}34+}
35+#endif
Mrandom/stateless_uniform/CMakeLists.txt+2-2
@@ -9,7 +9,7 @@
9# ---------------------------------------------------------------------------------------------------------9# ---------------------------------------------------------------------------------------------------------
10 10 
11# 设置算子定义时支持的芯片类型11# 设置算子定义时支持的芯片类型
12-set(SUPPORT_COMPUTE_UNIT "ascend950" "ascend350")12+set(SUPPORT_COMPUTE_UNIT "ascend950" "ascend350" "ascend960dt" "ascend960pr")
13# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译13# 设置每种芯片类型对应的tiling文件目录，即采用op_host目录下哪个文件夹下的tiling文件编译
14-set(SUPPORT_TILING_DIR "arch35" "arch35")14+set(SUPPORT_TILING_DIR "arch35" "arch35" "arch35" "arch35")
15add_all_modules_sources(OPTYPE stateless_uniform ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DEPENDENCIES random_common DISABLE_IN_OPP TRUE)15add_all_modules_sources(OPTYPE stateless_uniform ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DEPENDENCIES random_common DISABLE_IN_OPP TRUE)
Arandom/stateless_uniform/op_host/config/ascend960dt/stateless_uniform_simplified_key.ini+3-0
@@ -0,0 +1,3 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值
2+[StatelessUniform]
3+simplified_key_mode=0
Arandom/stateless_uniform/op_host/config/ascend960pr/stateless_uniform_simplified_key.ini+3-0
@@ -0,0 +1,3 @@
1+; 该文件主要影响 opc 工具 编译二进制kernel时， --simplified_key_mode 选项中填写的值
2+[StatelessUniform]
3+simplified_key_mode=0
Mrandom/stateless_uniform/op_host/stateless_uniform_def.cpp+2-0
@@ -84,6 +84,8 @@ public:
84 .PrecisionReduceFlag(true);84 .PrecisionReduceFlag(true);
85 this->AICore().AddConfig("ascend950");85 this->AICore().AddConfig("ascend950");
86 this->AICore().AddConfig("ascend350");86 this->AICore().AddConfig("ascend350");
87+ this->AICore().AddConfig("ascend960dt");
88+ this->AICore().AddConfig("ascend960pr");
87 }89 }
88};90};
89 91 
Arandom/stateless_uniform/tests/ut/op_kernel/stateless_uniform_tiling.h+57-0
@@ -0,0 +1,57 @@
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+#ifndef STATELESS_UNIFORM_TILING_H
12+#define STATELESS_UNIFORM_TILING_H
13+ 
14+#include <cstdint>
15+#include <cstring>
16+ 
17+#include "../../../../random_common/op_kernel/arch35/random_unified_tiling_data_arch35.h"
18+#include "kernel_tiling/kernel_tiling.h"
19+ 
20+static inline unsigned long long __mul_i32toi64(unsigned int lhs, unsigned int rhs)
21+{
22+ return static_cast<unsigned long long>(lhs) * rhs;
23+}
24+ 
25+#define __aicore__
26+#ifdef __NPU_TILING__
27+inline[aicore] void InitTilingData(const __gm__ uint8_t* tiling, RandomUnifiedSimtTilingDataStruct* constData)
28+{
29+ const __gm__ uint32_t* src = reinterpret_cast<const __gm__ uint32_t*>(tiling);
30+ uint32_t* dst = reinterpret_cast<uint32_t*>(constData);
31+ for (size_t i = 0; i < sizeof(RandomUnifiedSimtTilingDataStruct) / sizeof(uint32_t); ++i) {
32+ *(dst + i) = *(src + i);
33+ }
34+}
35+#else
36+inline void InitTilingData(uint8_t* tiling, RandomUnifiedSimtTilingDataStruct* constData)
37+{
38+ std::memcpy(constData, tiling, sizeof(RandomUnifiedSimtTilingDataStruct));
39+}
40+#endif // __NPU_TILING__
41+ 
42+#define CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
43+ __ubuf__ tilingStruct* tilingDataPointer = reinterpret_cast<__ubuf__ tilingStruct*>( \
44+ reinterpret_cast<__ubuf__ uint8_t*>(tilingPointer));
45+ 
46+#define INIT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer) \
47+ CONVERT_TILING_DATA(tilingStruct, tilingDataPointer, tilingPointer);
48+ 
49+#define GET_TILING_DATA_WITH_STRUCT(tilingStruct, tilingData, tilingArg) \
50+ tilingStruct tilingData; \
51+ InitTilingData(tilingArg, &tilingData)
52+ 
53+#define GET_TILING_DATA(tilingData, tilingArg) \
54+ RandomUnifiedSimtTilingDataStruct tilingData; \
55+ InitTilingData(tilingArg, &tilingData)
56+ 
57+#endif // STATELESS_UNIFORM_TILING_H
Mscripts/kernel/binary_config/ascendc_config.json+3-3
@@ -434,7 +434,7 @@
434 {"name":"MatmulReduceScatter", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false, "compile_options": {"ascend910_93":["-DAICORE_EXCEPTION_RESTART"]}},434 {"name":"MatmulReduceScatter", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false, "compile_options": {"ascend910_93":["-DAICORE_EXCEPTION_RESTART"]}},
435 {"name":"MatmulReduceScatterV2", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},435 {"name":"MatmulReduceScatterV2", "compute_units": ["ascend950"], "auto_sync" : false, "impl_mode" : ""},
436 {"name":"MaskedSoftmaxWithRelPosBias", "compute_units": ["ascend910b", "ascend310p"], "auto_sync" : false},436 {"name":"MaskedSoftmaxWithRelPosBias", "compute_units": ["ascend910b", "ascend310p"], "auto_sync" : false},
437- {"name":"StatelessRandomNormalV2", "compute_units": ["ascend950", "ascend350"], "auto_sync" : true, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},437+ {"name":"StatelessRandomNormalV2", "compute_units": ["ascend950", "ascend350", "ascend960dt", "ascend960pr"], "auto_sync" : true, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend960dt": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend960pr": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
438 {"name":"MaskedSelectV3", "compute_units": ["ascend910b", "ascend910_93", "ascend950", "kirinx90", "kirin9030", "ascend350"], "auto_sync" : false, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},438 {"name":"MaskedSelectV3", "compute_units": ["ascend910b", "ascend910_93", "ascend950", "kirinx90", "kirin9030", "ascend350"], "auto_sync" : false, "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
439 {"name":"Sinkhorn", "compute_units": ["ascend910b", "ascend910_93", "ascend950"], "auto_sync" : true},439 {"name":"Sinkhorn", "compute_units": ["ascend910b", "ascend910_93", "ascend950"], "auto_sync" : true},
440 {"name":"TomeMerge", "compute_units": ["ascend310p", "ascend910b", "ascend910"], "auto_sync" : true},440 {"name":"TomeMerge", "compute_units": ["ascend310p", "ascend910b", "ascend910"], "auto_sync" : true},
@@ -657,11 +657,11 @@
657 {"name":"RightShift", "compute_units": ["ascend950", "ascend350"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},657 {"name":"RightShift", "compute_units": ["ascend950", "ascend350"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
658 {"name":"NanToNum", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},658 {"name":"NanToNum", "compute_units": ["ascend950"], "auto_sync": false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
659 {"name":"StatelessRandomUniformV3", "compute_units": ["ascend950", "ascend350"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},659 {"name":"StatelessRandomUniformV3", "compute_units": ["ascend950", "ascend350"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
660- {"name":"StatelessRandomNormalV3", "compute_units": ["ascend950", "ascend350"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},660+ {"name":"StatelessRandomNormalV3", "compute_units": ["ascend950", "ascend350", "ascend960dt", "ascend960pr"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend960dt": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend960pr": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
661 {"name":"StatelessNormal", "compute_units": ["ascend950", "ascend350", "ascend960dt", "ascend960pr"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend960dt": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend960pr": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},661 {"name":"StatelessNormal", "compute_units": ["ascend950", "ascend350", "ascend960dt", "ascend960pr"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend960dt": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend960pr": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
662 {"name":"StatelessTruncatedNormalV2", "compute_units": ["ascend950", "ascend350"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},662 {"name":"StatelessTruncatedNormalV2", "compute_units": ["ascend950", "ascend350"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
663 {"name":"Adds", "compute_units": ["ascend950", "ascend350", "ascend960dt"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"],"ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend960dt": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},663 {"name":"Adds", "compute_units": ["ascend950", "ascend350", "ascend960dt"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"],"ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend960dt": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
664- {"name":"StatelessUniform", "compute_units": ["ascend950", "ascend350"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},664+ {"name":"StatelessUniform", "compute_units": ["ascend950", "ascend350", "ascend960dt", "ascend960pr"], "auto_sync" : false, "impl_mode": "", "compile_options": {"ascend950": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend350": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend960dt": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"], "ascend960pr": ["-mllvm -cce-aicore-dcci-before-kernel-end=false"]}},
665 {"name":"RadixTopK", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false},665 {"name":"RadixTopK", "compute_units": ["ascend910b", "ascend910_93"], "auto_sync" : false},
666 {"name":"Rint", "compute_units": ["ascend950", "ascend350"], "auto_sync" : false},666 {"name":"Rint", "compute_units": ["ascend950", "ascend350"], "auto_sync" : false},
667 {"name":"Tan", "compute_units": ["mc62"], "compile_options": {"mc62": ["--cce-no-dcache-flush"]}},667 {"name":"Tan", "compute_units": ["mc62"], "compile_options": {"mc62": ["--cce-no-dcache-flush"]}},