已合并
support RandomUniformIntV2 #957
xuejinghui创建于 1月27日
support RandomUniformIntV2 #957
已合并
xuejinghui创建于 1月27日
已删除 :master合入到cann/ops-mathmaster
23 个文件变更+2244-0
Mdocs/zh/op_list.md+10-0
@@ -2134,6 +2134,16 @@
2134 <td>AI Core</td>2134 <td>AI Core</td>
2135 <td>按照概率p随机将输入中的元素置零,并将输出按照1/(1-p)的比例缩放。</td>2135 <td>按照概率p随机将输入中的元素置零,并将输出按照1/(1-p)的比例缩放。</td>
2136 </tr>2136 </tr>
2137+ <tr>
2138+ <td>random</td>
2139+ <td><a href="../../random/uniform_int_v2/README.md">uniform_int_v2</a></td>
2140+ <td>√</td>
2141+ <td>√</td>
2142+ <td>√</td>
2143+ <td>√</td>
2144+ <td>AI Core</td>
2145+ <td>生成均匀分布的随机整数数列。</td>
2146+ </tr>
2137 <tr>2147 <tr>
2138 <td>random</td>2148 <td>random</td>
2139 <td><a href="../../random/stateless_bernoulli/README.md">stateless_bernoulli</a></td>2149 <td><a href="../../random/stateless_bernoulli/README.md">stateless_bernoulli</a></td>
Arandom/random_uniform_int_v2/CMakeLists.txt+15-0
@@ -0,0 +1,15 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# 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+set(SUPPORT_COMPUTE_UNIT "ascend950")
13+# 设置每种芯片类型对应的tiling文件目录,即采用op_host目录下哪个文件夹下的tiling文件编译
14+set(SUPPORT_TILING_DIR "arch35")
15+add_all_modules_sources(OPTYPE random_uniform_int_v2 ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR})
Arandom/random_uniform_int_v2/README.md+102-0
@@ -0,0 +1,102 @@
1+# RandomUniformIntV2
C
Cchenjiao1月30日

交付件缺少docs/zh/op_list.md

likedislike
xuejinghui
1月30日 评论:
xuejinghui
1月30日 评论:
2+ 
3+## 产品支持情况
4+ 
5+| 产品 | 是否支持 |
6+| :----------------------------------------------------------------- | :------: |
7+| Ascend 950PR/Ascend 950DT | √ |
8+| Atlas A3 训练系列产品/Atlas A3 推理系列产品 | × |
9+| Atlas A2 训练系列产品/Atlas A2 推理系列产品 | × |
C
Cchenjiao1月30日
已过期

A2产品不对

likedislike
xuejinghui
1月30日 评论:
xuejinghui
1月30日 评论:
10+| Atlas 200I/500 A2 推理产品 | × |
11+| Atlas 推理系列产品 | × |
12+| Atlas 训练系列产品 | × |
13+ 
14+## 功能说明
15+ 
16+- 算子功能:返回均匀分布的随机整数数列。
17+ 
18+## 参数说明
19+ 
20+<table style="undefined;table-layout: fixed; width: 1005px"><colgroup>
21+ <col style="width: 140px">
22+ <col style="width: 140px">
23+ <col style="width: 180px">
24+ <col style="width: 213px">
25+ <col style="width: 100px">
26+ </colgroup>
27+ <thead>
28+ <tr>
29+ <th>参数名</th>
30+ <th>输入/输出/属性</th>
31+ <th>描述</th>
32+ <th>数据类型</th>
33+ <th>数据格式</th>
34+ </tr></thead>
35+ <tbody>
36+ <tr>
37+ <td>shape</td>
38+ <td>输入</td>
39+ <td>输出张量的形状。</td>
40+ <td>INT64、INT32</td>
41+ <td>ND</td>
42+ </tr>
43+ <tr>
44+ <td>min</td>
45+ <td>输入</td>
46+ <td>生成的整数的包含下界。</td>
47+ <td>INT64、INT32</td>
48+ <td>ND</td>
49+ </tr>
50+ <tr>
51+ <td>max</td>
52+ <td>输入</td>
53+ <td>生成的整数的排除上界。</td>
54+ <td>INT64、INT32</td>
55+ <td>ND</td>
56+ </tr>
57+ <tr>
58+ <td>offset</td>
59+ <td>输入</td>
60+ <td>偏移值。</td>
61+ <td>INT64</td>
62+ <td>ND</td>
63+ </tr>
64+ <tr>
65+ <td>seed</td>
66+ <td>属性</td>
67+ <td>随机数种子。</td>
68+ <td>INT</td>
69+ <td>ND</td>
70+ </tr>
71+ <tr>
72+ <td>seed2</td>
73+ <td>属性</td>
74+ <td>随机数种子。</td>
75+ <td>INT</td>
76+ <td>ND</td>
77+ </tr>
78+ <tr>
79+ <td>y</td>
80+ <td>输出</td>
81+ <td>生成的随机数序列。</td>
82+ <td>INT64、INT32</td>
83+ <td>ND</td>
84+ </tr>
85+ <tr>
86+ <td>offset</td>
87+ <td>输出</td>
88+ <td>偏移值。</td>
89+ <td>INT64</td>
90+ <td>ND</td>
91+ </tr>
92+ </tbody></table>
93+ 
94+## 约束说明
95+ 
96+
97+ 
98+## 调用说明
99+ 
100+| 调用方式 | 样例代码 | 说明 |
101+| ---------- | ------------------------------------------------------------------------------ | ---------------------------------------------------------------------------------------- |
102+| 图模式调用 | [test_geir_random_uniform_int_v2](./examples/test_geir_random_uniform_int_v2.cpp) | 通过[算子IR](./op_graph/random_uniform_int_v2_proto.h)构图方式调用RandomUniformIntV2算子。 |
Arandom/random_uniform_int_v2/examples/test_geir_random_uniform_int_v2.cpp+336-0
@@ -0,0 +1,336 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <fstream>
13+#include <string.h>
14+#include <stdint.h>
15+#include <vector>
16+#include <string>
17+#include <map>
18+#include "assert.h"
19+ 
20+#include "graph.h"
21+#include "types.h"
22+#include "tensor.h"
23+#include "ge_error_codes.h"
24+#include "ge_api_types.h"
25+#include "ge_api.h"
26+#include "array_ops.h"
27+#include "ge_ir_build.h"
28+ 
29+#include "experiment_ops.h"
30+#include "nn_other.h"
31+#include "../op_graph/random_uniform_int_v2_proto.h"
32+ 
33+#define FAILED -1
34+#define SUCCESS 0
35+ 
36+using namespace ge;
37+using std::map;
38+using std::string;
39+using std::vector;
40+#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape, value) \
41+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
42+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
43+ TensorDesc placeholder##intputIndex##_desc = \
44+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
45+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
46+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
47+ Tensor tensor_placeholder##intputIndex; \
48+ ret = GenOnesDataFloat32(placeholder##intputIndex##_shape, \
49+ tensor_placeholder##intputIndex, \
50+ placeholder##intputIndex##_desc, \
51+ value); \
52+ if (ret != SUCCESS) { \
53+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
54+ return FAILED; \
55+ } \
56+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
57+ input.push_back(tensor_placeholder##intputIndex); \
58+ graph.AddOp(placeholder##intputIndex); \
59+ add1.set_input_##intputName(placeholder##intputIndex); \
60+ inputs.push_back(placeholder##intputIndex);
61+ 
62+#define ADD_INT_INPUT(intputIndex, intputName, intputDtype, inputShape, value) \
63+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
64+ auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
65+ TensorDesc placeholder##intputIndex##_desc = \
66+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
67+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
68+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
69+ Tensor tensor_placeholder##intputIndex; \
70+ ret = GenOnesData(placeholder##intputIndex##_shape, \
71+ tensor_placeholder##intputIndex, \
72+ placeholder##intputIndex##_desc, \
73+ intputDtype, \
74+ value); \
75+ if (ret != SUCCESS) { \
76+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
77+ return FAILED; \
78+ } \
79+ placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \
80+ input.push_back(tensor_placeholder##intputIndex); \
81+ graph.AddOp(placeholder##intputIndex); \
82+ add1.set_input_##intputName(placeholder##intputIndex); \
83+ inputs.push_back(placeholder##intputIndex);
84+ 
85+#define ADD_INPUT_ATTR(attrName, attrValue) \
86+ add1.set_attr_##attrName(attrValue);
87+ 
88+#define ADD_OUTPUT(outputIndex, outputName, outputDtype, outputShape) \
89+ TensorDesc outputName##outputIndex##_desc = \
90+ TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \
91+ add1.update_output_desc_##outputName(outputName##outputIndex##_desc);
92+ 
93+#define ADD_CONST_INPUT(intputIndex, intputName, intputDtype, inputShape) \
94+ vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
95+ auto placeholder##intputIndex = op::Const("placeholder" + intputIndex); \
96+ TensorDesc placeholder##intputIndex##_desc = \
97+ TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \
98+ placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \
99+ placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \
100+ Tensor tensor_placeholder##intputIndex; \
101+ ret = GenOnesData(placeholder##intputIndex##_shape, \
102+ tensor_placeholder##intputIndex, \
103+ placeholder##intputIndex##_desc, \
104+ intputDtype, \
105+ 1); \
106+ if (ret != SUCCESS) { \
107+ printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \
108+ return FAILED; \
109+ } \
110+ placeholder##intputIndex.SetAttr("value", tensor_placeholder##intputIndex); \
111+ placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \
112+ graph.AddOp(placeholder##intputIndex); \
113+ add1.set_input_##intputName(placeholder##intputIndex); \
114+ add1.update_input_desc_##intputName(placeholder##intputIndex##_desc); \
115+ inputs.push_back(placeholder##intputIndex);
116+ 
117+#define LOG_PRINT(message, ...) \
118+ do { \
119+ printf(message, ##__VA_ARGS__); \
120+ } while (0)
121+ 
122+string GetTime()
123+{
124+ time_t timep;
125+ time(&timep);
126+ char tmp[64];
127+ strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep));
128+ return tmp;
129+}
130+ 
131+uint32_t GetDataTypeSize(DataType dt)
132+{
133+ uint32_t dilation = 1;
134+ uint32_t oneByte = 1;
135+ uint32_t twoByte = 2;
136+ uint32_t fourByte = 4;
137+ uint32_t eightByte = 8;
138+ 
139+ if (dt == ge::DT_INT16) {
140+ dilation = twoByte;
141+ } else if (dt == ge::DT_UINT16) {
142+ dilation = twoByte;
143+ } else if (dt == ge::DT_INT32) {
144+ dilation = fourByte;
145+ } else if (dt == ge::DT_UINT32) {
146+ dilation = fourByte;
147+ } else if (dt == ge::DT_INT64) {
148+ dilation = eightByte;
149+ } else if (dt == ge::DT_UINT64) {
150+ dilation = eightByte;
151+ } else if (dt == ge::DT_INT8) {
152+ dilation = oneByte;
153+ }
154+ return dilation;
155+}
156+ 
157+int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor &input_tensor, TensorDesc &input_tensor_desc, float value)
158+{
159+ input_tensor_desc.SetRealDimCnt(shapes.size());
160+ size_t size = 1;
161+ for (uint32_t i = 0; i < shapes.size(); i++) {
162+ size *= shapes[i];
163+ }
164+ uint32_t byteSizeFloat32 = 4;
165+ uint32_t data_len = size * byteSizeFloat32;
166+ float *pData = new (std::nothrow) float[size];
167+ 
168+ for (size_t i = 0; i < size; ++i) {
169+ *(pData + i) = value;
170+ }
171+ input_tensor = Tensor(input_tensor_desc, (uint8_t *)pData, data_len);
172+ return SUCCESS;
173+}
174+ 
175+int32_t GenOnesData(
176+ vector<int64_t> shapes, Tensor &input_tensor, TensorDesc &input_tensor_desc, DataType data_type, int value)
177+{
178+ input_tensor_desc.SetRealDimCnt(shapes.size());
179+ size_t size = 1;
180+ for (uint32_t i = 0; i < shapes.size(); i++) {
181+ size *= shapes[i];
182+ }
183+ uint32_t data_len = size * GetDataTypeSize(data_type);
184+ int32_t *pData = new (std::nothrow) int32_t[data_len];
185+ for (uint32_t i = 0; i < size; ++i) {
186+ *(pData + i) = value;
187+ }
188+ input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t *>(pData), data_len);
189+ return SUCCESS;
190+}
191+ 
192+int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t *inputData)
193+{
194+ FILE *fp;
195+ fp = fopen(bin_file.c_str(), "w");
196+ fwrite(inputData, sizeof(uint8_t), data_size, fp);
197+ fclose(fp);
198+ return SUCCESS;
199+}
200+ 
201+int CreateOppInGraph(DataType inDtype, std::vector<ge::Tensor> &input, std::vector<Operator> &inputs,
202+ std::vector<Operator> &outputs, Graph &graph)
203+{
204+ Status ret = SUCCESS;
205+ // 自定义代码:添加单算子定义到图中
206+ auto add1 = op::RandomUniformIntV2("add1");
207+ std::vector<int64_t> xShape = {2};
208+ std::vector<int64_t> minShape = {1};
209+ std::vector<int64_t> maxShape = {1};
210+ std::vector<int64_t> offsetShape = {1};
211+ std::vector<int64_t> outShape = {1, 1};
212+ ADD_CONST_INPUT(1, shape, inDtype, xShape);
213+ ADD_INT_INPUT(2, min, inDtype, minShape, 0);
214+ ADD_INT_INPUT(3, max, inDtype, maxShape, 10);
215+ ADD_INT_INPUT(4, offset, inDtype, offsetShape, 0);
216+ 
217+ ADD_INPUT_ATTR(seed, 10);
218+ ADD_INPUT_ATTR(seed2, 5);
219+
220+ ADD_OUTPUT(1, y, ge::DT_INT64, outShape);
221+ ADD_OUTPUT(2, offset, ge::DT_INT64, offsetShape);
222+ 
223+ outputs.push_back(add1);
224+ // 添加完毕
225+ return SUCCESS;
226+}
227+ 
228+int main(int argc, char *argv[])
229+{
230+ const char *graph_name = "tc_ge_irrun_test";
231+ Graph graph(graph_name);
232+ std::vector<ge::Tensor> input;
233+ 
234+ printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str());
235+ std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}};
236+ Status ret = ge::GEInitialize(global_options);
237+ if (ret != SUCCESS) {
238+ printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str());
239+ return FAILED;
240+ }
241+ printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str());
242+ 
243+ std::vector<Operator> inputs{};
244+ std::vector<Operator> outputs{};
245+ 
246+ std::cout << argv[1] << std::endl;
247+ char *endptr;
248+ 
249+ DataType inDtype = DT_INT64;
250+ std::cout << inDtype << std::endl;
251+ 
252+ ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph);
253+ if (ret != SUCCESS) {
254+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
255+ return FAILED;
256+ }
257+ 
258+ if (!inputs.empty() && !outputs.empty()) {
259+ graph.SetInputs(inputs).SetOutputs(outputs);
260+ }
261+ 
262+ std::map<AscendString, AscendString> build_options = {
263+ 
264+ };
265+ printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str());
266+ ge::Session *session = new Session(build_options);
267+ 
268+ if (session == nullptr) {
269+ printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str());
270+ return FAILED;
271+ }
272+ printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str());
273+ printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str());
274+ 
275+ std::map<AscendString, AscendString> graph_options = {
276+ 
277+ };
278+ uint32_t graph_id = 0;
279+ ret = session->AddGraph(graph_id, graph, graph_options);
280+ 
281+ printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str());
282+ printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str());
283+ std::string file_path = "./dump";
284+ aclgrphDumpGraph(graph, file_path.c_str(), file_path.length());
285+ printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str());
286+ std::vector<ge::Tensor> output;
287+ ret = session->RunGraph(graph_id, input, output);
288+ if (ret != SUCCESS) {
289+ printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str());
290+ delete session;
291+ GEFinalize();
292+ return FAILED;
293+ }
294+ printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str());
295+ 
296+ int input_num = input.size();
297+ for (int i = 0; i < input_num; i++) {
298+ std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl;
299+ string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin";
300+ uint8_t *input_data_i = input[i].GetData();
301+ int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize();
302+ std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl;
303+ uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType());
304+ WriteDataToFile((const char *)input_file.c_str(), data_size, input_data_i);
305+ }
306+ 
307+ int output_num = output.size();
308+ for (int i = 0; i < output_num; i++) {
309+ std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl;
310+ string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin";
311+ uint8_t *output_data_i = output[i].GetData();
312+ int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize();
313+ std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl;
314+ uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType());
315+ WriteDataToFile((const char *)output_file.c_str(), data_size, output_data_i);
316+ int64_t *resultData = (int64_t*)output_data_i;
317+ for (int64_t j = 0; j < output_shape; j++) {
318+ LOG_PRINT("result[%ld] is: %f\n", j, resultData[j]);
319+ }
320+ }
321+ 
322+ ge::AscendString error_msg = ge::GEGetErrorMsgV2();
323+ std::string error_str(error_msg.GetString());
324+ std::cout << "Error message: " << error_str << std::endl;
325+ ge::AscendString warning_msg = ge::GEGetWarningMsgV2();
326+ std::string warning_str(warning_msg.GetString());
327+ std::cout << "Warning message: " << warning_str << std::endl;
328+ printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
329+ ret = ge::GEFinalize();
330+ if (ret != SUCCESS) {
331+ printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());
332+ return FAILED;
333+ }
334+ printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str());
335+ return SUCCESS;
336+}
Arandom/random_uniform_int_v2/op_graph/CMakeLists.txt+18-0
@@ -0,0 +1,18 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+foreach(SUB_DIR ${CURRENT_DIRS})
13+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14+ add_subdirectory(${SUB_DIR})
15+ endif()
16+endforeach()
17+ 
18+add_graph_plugin_sources()
Arandom/random_uniform_int_v2/op_graph/fusion_pass/.gitkeep+0-0
The file is empty
Arandom/random_uniform_int_v2/op_graph/random_uniform_int_v2_graph_infer.cpp+32-0
@@ -0,0 +1,32 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+/*!
11+ * \file random_uniform_int_v2_graph_infer.cpp
12+ * \brief random_uniform_int_v2 operator graph infer resource
13+ */
14+ 
15+#include "log/log.h"
16+#include "register/op_impl_registry.h"
17+ 
18+using namespace ge;
19+namespace ops {
20+ 
21+static graphStatus RandomUniformIntV2InferDataType(gert::InferDataTypeContext* context)
22+{
23+ OP_LOGD(context->GetNodeName(), "Begin to do RandomUniformIntV2InferDataType");
24+ // 实现推导逻辑
25+ DataType inputDataType = context->GetInputDataType(1);
26+ context->SetOutputDataType(0, inputDataType);
27+ context->SetOutputDataType(1, ge::DT_INT64);
28+ return ge::GRAPH_SUCCESS;
29+}
30+ 
31+IMPL_OP(RandomUniformIntV2).InferDataType(RandomUniformIntV2InferDataType);
CANN-robot
CANN-robotCANN-robot1月27日

代码结构与可维护性: 文件末尾缺少换行符(No newline at end of file),这不符合大多数编码规范。虽然这不会导致功能问题,但可能在某些工具(如diff、编译器)中产生警告,并且不符合良好的代码风格。

问题类型: 代码结构与可维护性 文件路径: random/random_uniform_int_v2/op_graph/random_uniform_int_v2_graph_infer.cpp 行号: 31 问题代码:

IMPL_OP(RandomUniformIntV2).InferDataType(RandomUniformIntV2InferDataType);
} // namespace ops

修改建议:

在文件末尾添加一个空行,确保以换行符结束。这符合POSIX标准,也是良好的编程实践。

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

likedislike
xuejinghui
1月29日 评论:
32+} // namespace ops
Arandom/random_uniform_int_v2/op_graph/random_uniform_int_v2_proto.h+47-0
@@ -0,0 +1,47 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * 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 RANDOM_UNIFORM_INT_V2_PROTO_H
12+#define RANDOM_UNIFORM_INT_V2_PROTO_H
13+ 
14+namespace ge {
15+/**
16+*@brief Outputs random values from a uniform distribution. \n
17+ 
18+*@par Inputs:
19+*Inputs include:
20+*@li shape: A 1-D Tensor. Must be one of the following types: int32, int64. The shape of the output tensor.
21+*@li min: A 1-D Tensor. Must be one of the following types: int32, int64.
22+*@li max: A 1-D Tensor. Must be one of the following types: int32, int64.
23+*@li offset: A 1-D Tensor, should be const data. Must be one of the following types: int64.
24+The value of offset should not be less than 0 and will be set to the default value 0 if it is negative. \n
CANN-robot
CANN-robotCANN-robot1月27日
已过期

代码结构与可维护性: 注释中存在拼写错误,'defualt'应为'default'。虽然这只是注释,但拼写错误会影响代码的可读性和专业性,特别是在技术文档中。

问题类型: 代码结构与可维护性 文件路径: random/random_uniform_int_v2/op_graph/random_uniform_int_v2_proto.h 行号: 24 问题代码:

The value of offset should not be less than 0 and will be set to the defualt value 0 if it is negative.

修改建议:

将注释中的'defualt'更正为'default'。

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

likedislike
xuejinghui
1月29日 评论:
25+ 
26+*@par Attributes:
27+*@li seed: An optional int. Defaults to 0. If either seed or seed2 are set to be non-zero,
28+the random number generator is seeded by the given seed. Otherwise, it is seeded by a random seed.
29+*@li seed2: An optional int. Defaults to 0 . A second seed to avoid seed collision. \n
30+ 
31+*@par Outputs:
32+*@li y: A Tensor. Has the same type as min. \n
33+*@li offset: A 1-D Tensor, should be const data. Must be one of the following types: int64. \n
CANN-robot
CANN-robotCANN-robot1月27日
已过期

代码结构与可维护性: 输出张量'offset'的注释描述与输入张量'offset'的注释描述不一致。输入描述为'should be const data',而输出描述中缺少此关键约束信息。这可能导致开发者误解输出张量的属性,认为它是可变的,从而引发潜在的错误使用。

问题类型: 代码结构与可维护性 文件路径: random/random_uniform_int_v2/op_graph/random_uniform_int_v2_proto.h 行号: 33 问题代码:

*@li offset: A 1-D Tensor, should be const data. Must be one of the following types: int64. \n

修改建议:

统一输入和输出张量'offset'的注释描述。建议在输出张量的注释中也明确说明其'should be const data'的属性,或者如果输出确实是可变的,则需要在输入注释中澄清其'const'特性仅适用于输入时刻。

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

likedislike
xuejinghui
1月29日 评论:
34+*/
35+REG_OP(RandomUniformIntV2)
36+ .INPUT(shape, TensorType({DT_INT32, DT_INT64}))
37+ .INPUT(min, TensorType({DT_INT32, DT_INT64}))
38+ .INPUT(max, TensorType({DT_INT32, DT_INT64}))
39+ .INPUT(offset, TensorType({DT_INT64}))
40+ .OUTPUT(y, TensorType({DT_INT32, DT_INT64}))
41+ .OUTPUT(offset, TensorType({DT_INT64}))
42+ .ATTR(seed, Int, 0)
43+ .ATTR(seed2, Int, 0)
44+ .OP_END_FACTORY_REG(RandomUniformIntV2)
45+}
46+ 
47+#endif
Arandom/random_uniform_int_v2/op_host/arch35/random_uniform_int_v2_tiling_arch35.cpp+369-0
@@ -0,0 +1,369 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * 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 random_uniform_int_v2_tiling_arch35.cpp
13+ * \brief
14+ */
15+ 
16+#include <random>
17+#include <iostream>
18+#include "platform/platform_infos_def.h"
19+#include "platform/platform_ascendc.h"
20+#include "op_common/op_host/util/platform_util.h"
21+#include "random_uniform_int_v2_tiling_arch35.h"
22+#include "../../../common/random_tiling_base.h"
23+#include "op_host/tiling_templates_registry.h"
24+#include "register/op_def_registry.h"
25+ 
26+namespace optiling {
27+ 
28+template <typename T>
29+ge::graphStatus RandomUniformIntV2Tiling::GetIntValue(const gert::Tensor *constTensor, gert::Shape &constShape)
30+{
31+ OP_LOGI(opName_, "RandomUniformIntV2Tiling::GetIntValue begin.");
32+ const T *constValue = constTensor->GetData<T>();
33+ OP_CHECK_NULL_WITH_CONTEXT(context_, constValue);
34+ const size_t constNum = constTensor->GetShapeSize();
35+ constShape.SetDimNum(0);
36+ for (size_t i = 0; i < constNum; ++i) {
37+ constShape.AppendDim(constValue[i]);
38+ }
39+ OP_LOGI(opName_, "RandomUniformIntV2Tiling::GetIntValue constNum=%zu", constNum);
40+ return ge::GRAPH_SUCCESS;
41+}
42+ 
43+ge::graphStatus RandomUniformIntV2Tiling::GetIntValueByDtype(const gert::Tensor *constTensor, gert::Shape &constShape,
44+ ge::DataType dType)
45+{
46+ ge::graphStatus ret = ge::GRAPH_SUCCESS;
47+ if (dType == ge::DataType::DT_INT32) {
48+ ret = GetIntValue<int32_t>(constTensor, constShape);
49+ } else {
50+ ret = GetIntValue<int64_t>(constTensor, constShape);
51+ }
52+ return ret;
53+}
54+ 
55+ge::graphStatus RandomUniformIntV2Tiling::GetMinAndMaxValue()
56+{
57+ OP_LOGI(opName_, "RandomUniformIntV2Tiling::GetMinAndMaxValue begin.");
58+ auto minDesc = context_->GetRequiredInputDesc(IN_MIN_IDX);
59+ OP_CHECK_NULL_WITH_CONTEXT(context_, minDesc);
60+ minDtype_ = minDesc->GetDataType();
61+ OP_CHECK_IF((minDtype_ != ge::DataType::DT_INT32) && (minDtype_ != ge::DataType::DT_INT64),
62+ OP_LOGE(opName_, "input min dtype should be int32, int64, but got %s.",
63+ Ops::Base::ToString(minDtype_).c_str()), return ge::GRAPH_FAILED);
64+ 
65+ auto minTensor = context_->GetRequiredInputTensor(IN_MIN_IDX);
66+ OP_CHECK_NULL_WITH_CONTEXT(context_, minTensor);
67+ auto minTensorSize = static_cast<int64_t>(minTensor->GetShapeSize());
68+ OP_CHECK_IF(minTensorSize != 1,
69+ OP_LOGE(opName_, "min data shape_size should be 1, but got %ld.", minTensorSize),
70+ return ge::GRAPH_FAILED);
71+ gert::Shape minShape;
72+ auto ret = GetIntValueByDtype(minTensor, minShape, minDtype_);
73+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS,
74+ OP_LOGE(opName_, "min GetIntValueByDtype failed."), return ge::GRAPH_FAILED);
75+ lo_ = static_cast<int64_t>(minShape.GetDim((0)));
76+ 
77+ auto maxDesc = context_->GetRequiredInputDesc(IN_MAX_IDX);
78+ OP_CHECK_NULL_WITH_CONTEXT(context_, maxDesc);
79+ auto maxDtype = maxDesc->GetDataType();
80+ OP_CHECK_IF(maxDtype != minDtype_,
81+ OP_LOGE(opName_, "input max dtype should have the same type as min, but got %s.",
82+ Ops::Base::ToString(maxDtype).c_str()), return ge::GRAPH_FAILED);
83+ 
84+ auto maxTensor = context_->GetRequiredInputTensor(IN_MAX_IDX);
85+ OP_CHECK_NULL_WITH_CONTEXT(context_, maxTensor);
86+ auto maxTensorSize = static_cast<int64_t>(maxTensor->GetShapeSize());
87+ OP_CHECK_IF(maxTensorSize != 1,
88+ OP_LOGE(opName_, "max data shape_size should be 1, but got %ld.", maxTensorSize),
89+ return ge::GRAPH_FAILED);
90+ gert::Shape maxShape;
91+ ret = GetIntValueByDtype(maxTensor, maxShape, maxDtype);
92+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS,
93+ OP_LOGE(opName_, "max GetIntValueByDtype failed."), return ge::GRAPH_FAILED);
94+ const int64_t maxTensorValue = static_cast<int64_t>(maxShape.GetDim((0)));
95+ OP_CHECK_IF(maxTensorValue <= lo_,
96+ OP_LOGE(opName_, "max should not be smaller or equal to min, but got max %ld, min %ld.",
97+ maxTensorValue, lo_),
98+ return ge::GRAPH_FAILED);
99+ range_ = static_cast<uint64_t>(maxTensorValue) - static_cast<uint64_t>(lo_);
100+ 
101+ OP_LOGI(opName_, "RandomUniformIntV2Tiling::GetMinAndMaxValue end.");
102+ return ge::GRAPH_SUCCESS;
103+}
104+ 
105+// 1、获取平台信息比如CoreNum、UB/L1/L0C资源大小
106+ge::graphStatus RandomUniformIntV2Tiling::GetPlatformInfo()
107+{
108+ OP_LOGI(opName_, "RandomUniformIntV2Tiling GetPlatformInfo.");
109+ auto compileInfo = static_cast<const RandomUniformIntV2CompileInfo*>(context_->GetCompileInfo());
110+ OP_CHECK_NULL_WITH_CONTEXT(context_, compileInfo);
111+ 
112+ totalCoreNum_ = static_cast<int64_t>(compileInfo->totalCoreNum);
113+ ubSize_ = compileInfo->ubSize;
114+ OP_CHECK_IF((ubSize_ <= 0), OP_LOGE(opName_, "ub size is invalid."), return ge::GRAPH_FAILED);
115+ OP_LOGI(opName_, "RandomUniformIntV2Tiling::GetPlatformInfo ubSize_=%d, totalCoreNum_=%d", ubSize_, totalCoreNum_);
116+ return ge::GRAPH_SUCCESS;
117+}
118+ 
119+// 2、获取INPUT/OUTPUT/ATTR信息
120+ge::graphStatus RandomUniformIntV2Tiling::GetShapeAttrsInfo()
121+{
122+ OP_LOGI(opName_, "RandomUniformIntV2Tiling::GetShapeAttrsInfo begin.");
123+ OP_CHECK_IF(GetInputInfo(),
124+ OP_LOGE(opName_, "GetInputInfo failed!"), return ge::GRAPH_FAILED);
125+
126+ OP_CHECK_IF(GetOutputInfo(),
127+ OP_LOGE(opName_, "GetOutputInfo failed!"), return ge::GRAPH_FAILED);
128+
129+ OP_CHECK_IF(shapeSize_ != outputSize_,
130+ OP_LOGE(opName_, "shape size: %ld is not equal to out size: %ld.", shapeSize_, outputSize_), return ge::GRAPH_FAILED);
131+ 
132+ OP_CHECK_IF(GetAttrInfo(),
133+ OP_LOGE(opName_, "GetAttrInfo failed!"), return ge::GRAPH_FAILED);
134+ OP_LOGI(opName_, "RandomUniformIntV2Tiling::GetShapeAttrsInfo end.");
135+ return ge::GRAPH_SUCCESS;
136+}
137+ 
138+ge::graphStatus RandomUniformIntV2Tiling::GetInputInfo()
139+{
140+ OP_LOGI(opName_, "RandomUniformIntV2Tiling::GetInputInfo begin.");
141+ auto shapeDesc = context_->GetRequiredInputDesc(IN_SHAPE_IDX);
142+ OP_CHECK_NULL_WITH_CONTEXT(context_, shapeDesc);
143+ auto shapeDtype = shapeDesc->GetDataType();
144+ OP_CHECK_IF((shapeDtype != ge::DataType::DT_INT32) && (shapeDtype != ge::DataType::DT_INT64),
145+ OP_LOGE(opName_, "input shape dtype should be int32, int64, but got %s.",
146+ Ops::Base::ToString(shapeDtype).c_str()), return ge::GRAPH_FAILED);
147+ 
148+ auto input1Shape = context_->GetInputShape(IN_SHAPE_IDX);
149+ OP_CHECK_NULL_WITH_CONTEXT(context_, input1Shape);
150+ uint32_t shapeDimNum = input1Shape->GetStorageShape().GetDimNum();
151+ OP_CHECK_IF(shapeDimNum != 1,
152+ OP_LOGE(opName_, "input shape is 1D tensor, but got %u.",
153+ shapeDimNum), return ge::GRAPH_FAILED);
154+ 
155+ auto shapeTensor = context_->GetRequiredInputTensor(IN_SHAPE_IDX);
156+ OP_CHECK_NULL_WITH_CONTEXT(context_, shapeTensor);
157+ gert::Shape constShape;
158+ auto ret = GetIntValueByDtype(shapeTensor, constShape, shapeDtype);
159+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS,
160+ OP_LOGE(opName_, "input shape GetIntValueByDtype failed."), return ge::GRAPH_FAILED);
161+ OP_LOGD(opName_, "RandomUniformIntV2Tiling::GetInputInfo get shapeTensor end.");
162+ 
163+ uint32_t shapeRank = constShape.GetDimNum();
164+ for (uint32_t idx = 0; idx < shapeRank; idx++) {
165+ shapeSize_ *= static_cast<int64_t>(constShape.GetDim(idx));
166+ }
167+ OP_CHECK_IF(shapeSize_ == 0,
168+ OP_LOGE(opName_, "input shape should not be empty tensor."), return ge::GRAPH_FAILED);
169+ 
170+ auto offsetDesc = context_->GetInputDesc(IN_OFFSET_IDX);
171+ OP_CHECK_NULL_WITH_CONTEXT(context_, offsetDesc);
172+ auto offsetDtype = offsetDesc->GetDataType();
173+ OP_CHECK_IF(offsetDtype != ge::DataType::DT_INT64,
174+ OP_LOGE(opName_, "input offset Dtype should be int64, but got %s.",
175+ Ops::Base::ToString(offsetDtype).c_str()), return ge::GRAPH_FAILED);
176+
177+ auto offsetTensor = context_->GetInputTensor(IN_OFFSET_IDX);
178+ OP_CHECK_NULL_WITH_CONTEXT(context_, offsetTensor);
179+ auto offsetTensorSize = static_cast<int64_t>(offsetTensor->GetShapeSize()); // 验证
180+ OP_CHECK_IF(offsetTensorSize != 1,
181+ OP_LOGE(opName_, "input offset shape_size should be 1, but got %ld.", offsetTensorSize), return ge::GRAPH_FAILED);
182+ 
183+ ret = GetMinAndMaxValue();
184+ OP_CHECK_IF(ret != ge::GRAPH_SUCCESS, OP_LOGE(opName_, "GetMinAndMaxValue failed."), return ge::GRAPH_FAILED);
185+ OP_LOGI(opName_, "RandomUniformIntV2Tiling::GetInputInfo end.");
186+ 
187+ return ge::GRAPH_SUCCESS;
188+}
189+ 
190+ge::graphStatus RandomUniformIntV2Tiling::GetOutputInfo()
191+{
192+ OP_LOGI(opName_, "RandomUniformIntV2Tiling::GetOutputInfo begin.");
193+ auto outDesc = context_->GetOutputDesc(OUTPUT_IDX_Y);
194+ OP_CHECK_NULL_WITH_CONTEXT(context_, outDesc);
195+ outDtype_ = outDesc->GetDataType();
196+ OP_CHECK_IF(outDtype_ != minDtype_,
197+ OP_LOGE(opName_, "out shape dtype should have the same type as min, but got %s.",
198+ Ops::Base::ToString(outDtype_).c_str()), return ge::GRAPH_FAILED);
199+ 
200+ auto outputShape = context_->GetOutputShape(OUTPUT_IDX_Y);
201+ OP_CHECK_NULL_WITH_CONTEXT(context_, outputShape);
202+ auto outTensor = outputShape->GetStorageShape();
203+ outputSize_ = outTensor.GetShapeSize();
204+ OP_CHECK_IF(outputSize_ == 0,
205+ OP_LOGE(opName_, "output shape_size should not be 0."), return ge::GRAPH_FAILED);
206+ OP_LOGI(opName_, "RandomUniformIntV2Tiling::GetOutputInfo end.");
207+ return ge::GRAPH_SUCCESS;
208+}
209+ 
210+ge::graphStatus RandomUniformIntV2Tiling::GetAttrInfo()
211+{
212+ OP_LOGI(opName_, "RandomUniformIntV2Tiling::GetAttrInfo begin.");
213+ auto attrs = context_->GetAttrs();
214+ OP_CHECK_NULL_WITH_CONTEXT(context_, attrs);
215+ const auto* seedAttr = attrs->GetAttrPointer<int64_t>(ATTR_SEED_IDX);
216+ OP_CHECK_NULL_WITH_CONTEXT(context_, seedAttr);
217+ const auto* seed2Attr = attrs->GetAttrPointer<int64_t>(ATTR_SEED2_IDX);
218+ OP_CHECK_NULL_WITH_CONTEXT(context_, seed2Attr);
219+
220+ seed_ = *seedAttr;
221+ seed2_ = *seed2Attr;
222+ if (seed_ == 0 && seed2_ == 0) {
223+ seed_ = static_cast<int64_t>(New64());
224+ seed2_ = static_cast<int64_t>(New64());
225+ }
226+ OP_LOGI(opName_, "RandomUniformIntV2Tiling seed_ is %ld, seed2_ is %ld", seed_, seed2_);
227+ return ge::GRAPH_SUCCESS;
228+}
229+ 
230+bool RandomUniformIntV2Tiling::IsCapable()
231+{
232+ return true;
233+}
234+ 
235+void RandomUniformIntV2Tiling::SetTilingData()
236+{
237+ RandomUniformIntV2TilingData4RegBase* tilingData = context_->GetTilingData<RandomUniformIntV2TilingData4RegBase>();
238+ tilingData->blockNum = blockNum_;
239+ tilingData->normalCoreProNum = normalCoreProNum_;
240+ tilingData->tailCoreProNum = tailCoreProNum_;
241+ tilingData->singleUbSize = singleUbSize_;
242+ tilingData->seed = seed_;
243+ tilingData->seed2 = seed2_;
244+ tilingData->outputSize = outputSize_;
245+ tilingData->range = range_;
246+ tilingData->lo = lo_;
247+}
248+ 
249+void RandomUniformIntV2Tiling::DoBlockTiling()
250+{
251+ outputDtypeSize_ = ge::GetSizeByDataType(outDtype_);
252+ if (outputDtypeSize_ == 0) {
253+ return;
254+ }
255+ 
256+ auto coreAlignFactor = CORE_ALIGN_SIZE / outputDtypeSize_;
257+ auto blockFactor = Ops::Base::CeilDiv(outputSize_, totalCoreNum_);
258+ auto blockAlignFactor = Ops::Base::CeilDiv(blockFactor, coreAlignFactor) * coreAlignFactor;
259+ auto minTilingSize = MIN_TILING_SIZE;
260+ normalCoreProNum_ = std::max(static_cast<uint32_t>(blockAlignFactor), minTilingSize);
261+ blockNum_ = Ops::Base::CeilDiv(outputSize_, normalCoreProNum_);
262+ tailCoreProNum_ = outputSize_ - normalCoreProNum_ * (blockNum_ - 1);
263+ return;
264+}
265+ 
266+void RandomUniformIntV2Tiling::UbTiling()
267+{
268+ // quarterUbSize: 2 for double buffer; coefVal for temp RNG, philox temp buff need uint32 to int32/int64
269+ int64_t coefVal = DOUBLE_BUFFER;
270+ auto quarterUbSize = (ubSize_ - DCACHE_SIZE) / (DOUBLE_BUFFER + coefVal);
271+ auto ubBlockSize = static_cast<int32_t>(Ops::Base::GetUbBlockSize(context_));
272+ auto alignFactor = ubBlockSize / outputDtypeSize_;
273+ singleUbSize_ = (quarterUbSize / outputDtypeSize_ / alignFactor) * alignFactor;
274+}
275+ 
276+// 3、计算数据切分TilingData
277+ge::graphStatus RandomUniformIntV2Tiling::DoOpTiling()
278+{
279+ OP_LOGD(opName_, "RandomUniformIntV2Tiling DoOpTiling.");
280+ DoBlockTiling();
281+ UbTiling();
282+ SetTilingData();
283+ return ge::GRAPH_SUCCESS;
284+}
285+ 
286+// 4、计算高阶API的TilingData
287+ge::graphStatus RandomUniformIntV2Tiling::DoLibApiTiling()
288+{
289+ return ge::GRAPH_SUCCESS;
290+}
291+ 
292+// 5、计算TilingKey
293+uint64_t RandomUniformIntV2Tiling::GetTilingKey() const
294+{
295+ uint64_t tilingKey = GET_TPL_TILING_KEY(RANDOM_UNIFORM_INT_V2_TPL);
296+ return tilingKey;
297+}
298+ 
299+// 6、计算Workspace 大小
300+ge::graphStatus RandomUniformIntV2Tiling::GetWorkspaceSize()
301+{
302+ workspaceSize_ = DEFAULT_WORKSPACE_SIZE;
303+ return ge::GRAPH_SUCCESS;
304+}
305+ 
306+// 7、保存Tiling数据
307+ge::graphStatus RandomUniformIntV2Tiling::PostTiling()
308+{
309+ auto workspaces = context_->GetWorkspaceSizes(1);
310+ OP_CHECK_NULL_WITH_CONTEXT(context_, workspaces);
311+ workspaces[0] = workspaceSize_;
312+ context_->SetBlockDim(blockNum_);
313+ context_->SetLocalMemorySize(ubSize_ - DCACHE_SIZE);
314+ context_->SetTilingKey(GetTilingKey());
315+ context_->SetScheduleMode(1);
316+ return ge::GRAPH_SUCCESS;
317+}
318+ 
319+void RandomUniformIntV2Tiling::DumpTilingInfo()
320+{
321+ std::ostringstream info;
322+ info << " ubSize: " << ubSize_;
323+ info << " totalCoreNum: " << totalCoreNum_;
324+ info << " blockNum: " << blockNum_;
325+ info << " normalCoreProNum: " << normalCoreProNum_;
326+ info << " tailCoreProNum: " << tailCoreProNum_;
327+ info << " singleUbSize: " << singleUbSize_;
328+ info << " seed: " << seed_;
329+ info << " seed2: " << seed2_;
330+ info << " outputSize: " << outputSize_;
331+ info << " range: " << range_;
332+ info << " lo: " << lo_;
333+ 
334+ OP_LOGI(opName_, "%s", info.str().c_str());
335+}
336+ 
337+static ge::graphStatus TilingPrepare4RandomUniformIntV2Tiling(gert::TilingParseContext* context)
338+{
339+ auto compileInfo = context->GetCompiledInfo<RandomUniformIntV2CompileInfo>();
340+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
341+ auto platformInfo = context->GetPlatformInfo();
342+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
343+ compileInfo->totalCoreNum = ascendcPlatform.GetCoreNumAiv();
344+ uint64_t ubSizePlatForm;
345+ ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSizePlatForm);
346+ compileInfo->ubSize = static_cast<int64_t>(ubSizePlatForm);
347+ OP_CHECK_IF(
348+ (compileInfo->totalCoreNum <= 0 || compileInfo->ubSize <= 0),
349+ OP_LOGE(
350+ context, "RandomUniformIntV2 GetHardwareInfo Failed, vectorCoreNum:%ld, ubSize:%ld.", compileInfo->totalCoreNum,
351+ compileInfo->ubSize),
352+ return ge::GRAPH_FAILED);
353+ OP_LOGD(context, "Get totalCoreNum:%d, ubSize:%ld", compileInfo->totalCoreNum, compileInfo->ubSize);
354+ return ge::GRAPH_SUCCESS;
355+}
356+ 
357+static ge::graphStatus TilingRandomUniformIntV2(gert::TilingContext* tilingContext)
358+{
359+ OP_LOGD(tilingContext, "Entering TilingRandomUniformIntV2");
360+ RandomUniformIntV2Tiling tilingObj(tilingContext);
361+ return tilingObj.DoTiling();
362+}
363+ 
364+IMPL_OP_OPTILING(RandomUniformIntV2)
365+ .Tiling(TilingRandomUniformIntV2)
366+ .TilingParse<RandomUniformIntV2CompileInfo>(TilingPrepare4RandomUniformIntV2Tiling)
367+ .TilingInputsDataDependency({IN_SHAPE_IDX, IN_MIN_IDX, IN_MAX_IDX});
368+ 
369+} // namespace optiling
Arandom/random_uniform_int_v2/op_host/arch35/random_uniform_int_v2_tiling_arch35.h+102-0
@@ -0,0 +1,102 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * 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 random_uniform_int_v2_tiling_arch35.h
13+ * \brief
14+ */
15+#ifndef RANDOM_UNIFORM_INT_V2_TILING_ARCH35_H
16+#define RANDOM_UNIFORM_INT_V2_TILING_ARCH35_H
17+ 
18+#include <string>
19+#include "op_host/tiling_base.h"
20+#include "util/math_util.h"
21+#include "../../op_kernel/arch35/random_uniform_int_v2_struct.h"
22+ 
23+namespace optiling {
24+ 
25+static constexpr uint16_t IN_SHAPE_IDX = 0;
26+static constexpr uint16_t IN_MIN_IDX = 1;
27+static constexpr uint16_t IN_MAX_IDX = 2;
28+static constexpr uint16_t IN_OFFSET_IDX = 3;
29+static constexpr uint16_t OUTPUT_IDX_Y = 0;
30+static constexpr uint16_t ATTR_SEED_IDX = 0;
31+static constexpr uint16_t ATTR_SEED2_IDX = 1;
32+static constexpr int64_t CORE_ALIGN_SIZE = 512;
33+static constexpr uint32_t MIN_TILING_SIZE = 256;
34+static constexpr uint64_t DEFAULT_WORKSPACE_SIZE = 0;
35+static constexpr int64_t DOUBLE_BUFFER = 2;
36+static constexpr int64_t DCACHE_SIZE = 32 * 1024;
CANN-robot
CANN-robotCANN-robot1月27日
已过期

代码风格-常量定义一致性: 常量DCACHE_SIZE的定义风格与其他常量不一致。其他常量使用static constexpr且类型在前(如static constexpr int64_t),而DCACHE_SIZE使用了constexpr static(static在constexpr后)。虽然这在语法上等效,但保持统一的代码风格有助于提高可读性。

问题类型: 代码风格-常量定义一致性 文件路径: random/random_uniform_int_v2/op_host/arch35/random_uniform_int_v2_tiling_arch35.h 行号: 36 问题代码:

constexpr static int64_t DCACHE_SIZE = 32 * 1024;

修改建议:

将DCACHE_SIZE的定义改为与其他常量一致的格式:static constexpr int64_t DCACHE_SIZE = 32 * 1024;

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

likedislike
xuejinghui
1月29日 评论:
37+ 
38+class RandomUniformIntV2Tiling : public Ops::Math::OpTiling::TilingBaseClass
39+{
40+public:
41+ explicit RandomUniformIntV2Tiling(gert::TilingContext* context)
42+ : TilingBaseClass(context), opName_(context->GetNodeName())
43+ {}
44+ 
45+protected:
46+ bool IsCapable() override;
47+ // 1、获取平台信息比如CoreNum、UB/L1/L0C资源大小
48+ ge::graphStatus GetPlatformInfo() override;
49+ // 2、获取INPUT/OUTPUT/ATTR信息
50+ ge::graphStatus GetShapeAttrsInfo() override;
51+ // 3、计算数据切分TilingData
52+ ge::graphStatus DoOpTiling() override;
53+ // 4、计算高阶API的TilingData
54+ ge::graphStatus DoLibApiTiling() override;
55+ // 5、计算TilingKey
56+ uint64_t GetTilingKey() const override;
57+ // 6、计算Workspace 大小
58+ ge::graphStatus GetWorkspaceSize() override;
59+ // 7、保存Tiling数据
60+ ge::graphStatus PostTiling() override;
61+ 
62+ void DumpTilingInfo() override;
63+ 
64+private:
65+ const std::string opName_;
66+ int64_t ubSize_{0};
67+ int64_t totalCoreNum_{0};
68+ int64_t shapeSize_{1};
69+ int64_t outputSize_{0};
70+ int64_t seed_{0};
71+ int64_t seed2_{0};
72+ ge::DataType outDtype_{ge::DT_INT64};
73+ ge::DataType minDtype_{ge::DT_INT64};
74+ 
75+ int64_t outputDtypeSize_{0};
76+ int64_t blockNum_{0};
77+ int64_t normalCoreProNum_{0};
78+ int64_t tailCoreProNum_{0};
79+ int64_t singleUbSize_{0};
80+ uint64_t range_{0};
81+ int64_t lo_{0};
82+ 
83+ void SetTilingData();
84+ template <typename T>
85+ ge::graphStatus GetIntValue(const gert::Tensor *constTensor, gert::Shape &constShape);
86+ ge::graphStatus GetIntValueByDtype(const gert::Tensor *constTensor, gert::Shape &constShape, ge::DataType dType);
87+ ge::graphStatus GetMinAndMaxValue();
88+ ge::graphStatus GetInputInfo();
89+ ge::graphStatus GetOutputInfo();
90+ ge::graphStatus GetAttrInfo();
91+ void DoBlockTiling();
92+ void UbTiling();
93+};
94+ 
95+struct RandomUniformIntV2CompileInfo {
96+ int64_t totalCoreNum = 0;
97+ int64_t ubSize = 0;
98+};
99+ 
100+} // namespace optiling
101+ 
102+#endif // RANDOM_UNIFORM_INT_V2_TILING_ARCH35_H
CANN-robot
CANN-robotCANN-robot1月27日

代码质量-文件结尾格式: 文件末尾缺少换行符(No newline at end of file)。虽然这不是功能性问题,但不符合POSIX标准,可能导致某些工具(如diff、cat)处理时出现格式问题,也可能被代码规范检查工具标记为警告。

问题类型: 代码质量-文件结尾格式 文件路径: random/random_uniform_int_v2/op_host/arch35/random_uniform_int_v2_tiling_arch35.h 行号: 102 问题代码:

#endif // RANDOM_UNIFORM_INT_V2_TILING_ARCH35_H\ No newline at end of file

修改建议:

在文件末尾添加一个空行,确保以换行符结束。

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

likedislike
xuejinghui
1月29日 评论:
Arandom/random_uniform_int_v2/op_host/config/ascend950/random_uniform_int_v2_binary.json+321-0
@@ -0,0 +1,321 @@
1+{
2+ "op_type": "RandomUniformIntV2",
3+ "op_list": [
4+ {
5+ "bin_filename": "RandomUniformIntV2_1ea4bdp73cd9df67072d9z7caec7c583",
6+ "inputs": [
7+ {
8+ "name": "shape",
9+ "index": 0,
10+ "dtype": "int64",
11+ "format": "ND",
12+ "shape": [
13+ -2
14+ ],
15+ "paramType": "required"
16+ },
17+ {
18+ "name": "min",
19+ "index": 1,
20+ "dtype": "int64",
21+ "format": "ND",
22+ "shape": [
23+ -2
24+ ],
25+ "paramType": "required"
26+ },
27+ {
28+ "name": "max",
29+ "index": 2,
30+ "dtype": "int64",
31+ "format": "ND",
32+ "shape": [
33+ -2
34+ ],
35+ "paramType": "required"
36+ },
37+ {
38+ "name": "offset",
39+ "index": 3,
40+ "dtype": "int64",
41+ "format": "ND",
42+ "shape": [
43+ -2
44+ ],
45+ "paramType": "required"
46+ }
47+ ],
48+ "outputs": [
49+ {
50+ "name": "y",
51+ "index": 0,
52+ "dtype": "int64",
53+ "format": "ND",
54+ "shape": [
55+ -2
56+ ],
57+ "paramType": "required"
58+ },
59+ {
60+ "name": "offset",
61+ "index": 1,
62+ "dtype": "int64",
63+ "format": "ND",
64+ "shape": [
65+ -2
66+ ],
67+ "paramType": "required"
68+ }
69+ ],
70+ "attrs": [
71+ {
72+ "name": "seed",
73+ "dtype": "int",
74+ "value": 0
75+ },
76+ {
77+ "name": "seed2",
78+ "dtype": "int",
79+ "value": 0
80+ }
81+ ]
82+ },
83+ {
84+ "bin_filename": "RandomUniformIntV2_1ea7bd0e3cd9ef67072d9z7amgc9c733",
85+ "inputs": [
86+ {
87+ "name": "shape",
88+ "index": 0,
89+ "dtype": "int32",
90+ "format": "ND",
91+ "shape": [
92+ -2
93+ ],
94+ "paramType": "required"
95+ },
96+ {
97+ "name": "min",
98+ "index": 1,
99+ "dtype": "int32",
100+ "format": "ND",
101+ "shape": [
102+ -2
103+ ],
104+ "paramType": "required"
105+ },
106+ {
107+ "name": "max",
108+ "index": 2,
109+ "dtype": "int32",
110+ "format": "ND",
111+ "shape": [
112+ -2
113+ ],
114+ "paramType": "required"
115+ },
116+ {
117+ "name": "offset",
118+ "index": 3,
119+ "dtype": "int64",
120+ "format": "ND",
121+ "shape": [
122+ -2
123+ ],
124+ "paramType": "required"
125+ }
126+ ],
127+ "outputs": [
128+ {
129+ "name": "y",
130+ "index": 0,
131+ "dtype": "int32",
132+ "format": "ND",
133+ "shape": [
134+ -2
135+ ],
136+ "paramType": "required"
137+ },
138+ {
139+ "name": "offset",
140+ "index": 1,
141+ "dtype": "int64",
142+ "format": "ND",
143+ "shape": [
144+ -2
145+ ],
146+ "paramType": "required"
147+ }
148+ ],
149+ "attrs": [
150+ {
151+ "name": "seed",
152+ "dtype": "int",
153+ "value": 0
154+ },
155+ {
156+ "name": "seed2",
157+ "dtype": "int",
158+ "value": 0
159+ }
160+ ]
161+ },
162+ {
163+ "bin_filename": "RandomUniformIntV2_1ea7bd0e3cd9de67072d9z7akyc9c583",
164+ "inputs": [
165+ {
166+ "name": "shape",
167+ "index": 0,
168+ "dtype": "int32",
169+ "format": "ND",
170+ "shape": [
171+ -2
172+ ],
173+ "paramType": "required"
174+ },
175+ {
176+ "name": "min",
177+ "index": 1,
178+ "dtype": "int64",
179+ "format": "ND",
180+ "shape": [
181+ -2
182+ ],
183+ "paramType": "required"
184+ },
185+ {
186+ "name": "max",
187+ "index": 2,
188+ "dtype": "int64",
189+ "format": "ND",
190+ "shape": [
191+ -2
192+ ],
193+ "paramType": "required"
194+ },
195+ {
196+ "name": "offset",
197+ "index": 3,
198+ "dtype": "int64",
199+ "format": "ND",
200+ "shape": [
201+ -2
202+ ],
203+ "paramType": "required"
204+ }
205+ ],
206+ "outputs": [
207+ {
208+ "name": "y",
209+ "index": 0,
210+ "dtype": "int64",
211+ "format": "ND",
212+ "shape": [
213+ -2
214+ ],
215+ "paramType": "required"
216+ },
217+ {
218+ "name": "offset",
219+ "index": 1,
220+ "dtype": "int64",
221+ "format": "ND",
222+ "shape": [
223+ -2
224+ ],
225+ "paramType": "required"
226+ }
227+ ],
228+ "attrs": [
229+ {
230+ "name": "seed",
231+ "dtype": "int",
232+ "value": 0
233+ },
234+ {
235+ "name": "seed2",
236+ "dtype": "int",
237+ "value": 0
238+ }
239+ ]
240+ },
241+ {
242+ "bin_filename": "RandomUniformIntV2_1ea7bd0e3cd9df67072d9b7aaec9c583",
243+ "inputs": [
244+ {
245+ "name": "shape",
246+ "index": 0,
247+ "dtype": "int64",
248+ "format": "ND",
249+ "shape": [
250+ -2
251+ ],
252+ "paramType": "required"
253+ },
254+ {
255+ "name": "min",
256+ "index": 1,
257+ "dtype": "int32",
258+ "format": "ND",
259+ "shape": [
260+ -2
261+ ],
262+ "paramType": "required"
263+ },
264+ {
265+ "name": "max",
266+ "index": 2,
267+ "dtype": "int32",
268+ "format": "ND",
269+ "shape": [
270+ -2
271+ ],
272+ "paramType": "required"
273+ },
274+ {
275+ "name": "offset",
276+ "index": 3,
277+ "dtype": "int64",
278+ "format": "ND",
279+ "shape": [
280+ -2
281+ ],
282+ "paramType": "required"
283+ }
284+ ],
285+ "outputs": [
286+ {
287+ "name": "y",
288+ "index": 0,
289+ "dtype": "int32",
290+ "format": "ND",
291+ "shape": [
292+ -2
293+ ],
294+ "paramType": "required"
295+ },
296+ {
297+ "name": "offset",
298+ "index": 1,
299+ "dtype": "int64",
300+ "format": "ND",
301+ "shape": [
302+ -2
303+ ],
304+ "paramType": "required"
305+ }
306+ ],
307+ "attrs": [
308+ {
309+ "name": "seed",
310+ "dtype": "int",
311+ "value": 0
312+ },
313+ {
314+ "name": "seed2",
315+ "dtype": "int",
316+ "value": 0
317+ }
318+ ]
319+ }
320+ ]
321+}
Arandom/random_uniform_int_v2/op_host/config/ascend950/random_uniform_int_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,ascendxx为可选mode,如果不同芯片有差异化要求时,需要配置;
6+; 1)如果没有配置:非ascendC算子继续按空处理,即opc编译命令中不添加 --simplified_key_mode 选项,AscendC算子按照 simplified_key_mode=0 处理
7+; 2)如果仅有default配置:各个版本按default配置
8+; 3)如果仅有某些平台的配置,没有default配置:对应平台的按照配置的值传递,非对应平台的:非AscendC算子继续按空处理,AscendC算子按照 simplified_key_mode=0 处理
9+; 4)如果default配置和平台配置都有:对应平台的使用平台的配置,非对应的平台的以default值配置。
10+; 5)对于自定义simplified key的情况,需要在binary_simplified_key_mode.ini 文件中显式配置为None,不传入 --simplified_key_mode 选项,由opc工具和FE框架自行判断使用何种模式
11+; 6)是否是AscendC算子,由 ops/build-in/tbe/op_info_cfg/parser/ascendc_config.json 中配置的算子名字和对于的平台决定
12+[RandomUniformIntV2]
13+default=0
Arandom/random_uniform_int_v2/op_host/random_uniform_int_v2_def.cpp+70-0
@@ -0,0 +1,70 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * 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 random_uniform_int_v2_def.cpp
13+ * \brief
14+ */
15+#include "register/op_def_registry.h"
16+ 
17+namespace ops {
18+class RandomUniformIntV2 : public OpDef {
19+public:
20+ explicit RandomUniformIntV2(const char* name) : OpDef(name)
21+ {
22+ this->Input("shape")
23+ .ParamType(REQUIRED)
24+ .DataType({ge::DT_INT64, ge::DT_INT64, ge::DT_INT32, ge::DT_INT32})
25+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
26+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
27+ .ValueDepend(OPTIONAL);
28+ this->Input("min")
29+ .ParamType(REQUIRED)
30+ .DataType({ge::DT_INT64, ge::DT_INT32, ge::DT_INT64, ge::DT_INT32})
31+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
32+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
33+ .ValueDepend(OPTIONAL);
34+ this->Input("max")
35+ .ParamType(REQUIRED)
36+ .DataType({ge::DT_INT64, ge::DT_INT32, ge::DT_INT64, ge::DT_INT32})
37+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
38+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
39+ .ValueDepend(OPTIONAL);
40+ this->Input("offset")
41+ .ParamType(REQUIRED)
42+ .DataType({ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64})
43+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
44+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
45+ this->Output("y")
46+ .ParamType(REQUIRED)
47+ .DataType({ge::DT_INT64, ge::DT_INT32, ge::DT_INT64, ge::DT_INT32})
48+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
49+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
50+ this->Output("offset")
51+ .ParamType(REQUIRED)
52+ .DataType({ge::DT_INT64, ge::DT_INT64, ge::DT_INT64, ge::DT_INT64})
53+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
54+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
55+ 
56+ this->Attr("seed").AttrType(OPTIONAL).Int(0);
57+ this->Attr("seed2").AttrType(OPTIONAL).Int(0);
58+ 
59+ OpAICoreConfig aicoreconfig;
60+ aicoreconfig.DynamicCompileStaticFlag(true)
61+ .DynamicFormatFlag(false)
62+ .DynamicRankSupportFlag(true)
63+ .DynamicShapeSupportFlag(true)
64+ .NeedCheckSupportFlag(false)
65+ .PrecisionReduceFlag(true);
66+ this->AICore().AddConfig("ascend950", aicoreconfig);
67+ }
68+};
69+OP_ADD(RandomUniformIntV2);
70+} // namespace ops
CANN-robot
CANN-robotCANN-robot1月27日

代码结构与可维护性: 文件末尾缺少换行符(No newline at end of file),这是代码规范问题。虽然不影响功能,但可能在某些工具(如diff、编译器)中产生警告,并且不符合通用的代码风格规范。

问题类型: 代码结构与可维护性 文件路径: random/random_uniform_int_v2/op_host/random_uniform_int_v2_def.cpp 行号: 70 问题代码:

} // namespace ops

修改建议:

在文件末尾添加一个空行,确保以换行符结束。

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

likedislike
xuejinghui
1月29日 评论:
Arandom/random_uniform_int_v2/op_host/random_uniform_int_v2_infershape.cpp+60-0
@@ -0,0 +1,60 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * 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 random_uniform_int_v2_infershape.cpp
13+ * \brief
14+ */
15+#include "log/log.h"
16+#include "register/op_impl_registry.h"
17+ 
18+using namespace ge;
19+namespace ops {
20+ 
21+static constexpr uint16_t IN_SHAPE_IDX = 0;
22+ 
23+template <typename T>
24+ge::graphStatus RandomUniformIntV2InferShapeImpl(const T *shape_dims, gert::Shape &outputShape,
25+ size_t shape_size) {
26+ outputShape.SetDimNum(shape_size);
27+ for (size_t i = 0U; i < shape_size; i++) {
28+ outputShape.SetDim(i, shape_dims[i]);
29+ }
30+ return ge::GRAPH_SUCCESS;
31+}
32+ 
33+static ge::graphStatus InferShapeForRandomUniformIntV2(gert::InferShapeContext *context) {
34+ auto x_shape_tensor = context->GetInputTensor(0);
35+ auto outputShape = context->GetOutputShape(0);
36+ auto const_shape = context->GetOutputShape(1);
37+ OP_CHECK_NULL_WITH_CONTEXT(context, x_shape_tensor);
38+ OP_CHECK_NULL_WITH_CONTEXT(context, outputShape);
39+ OP_CHECK_NULL_WITH_CONTEXT(context, const_shape);
40+ 
41+ const_shape->SetDimNum(1);
42+ const_shape->SetDim(0, 1);
43+ 
44+ auto x_shape_size = x_shape_tensor->GetShapeSize();
45+ if (x_shape_size < 0) {
46+ return ge::GRAPH_FAILED;
47+ }
48+ if (x_shape_tensor->GetDataType() == ge::DT_INT32) {
49+ auto xShapeData = x_shape_tensor->GetData<int32_t>();
50+ return RandomUniformIntV2InferShapeImpl<int32_t>(xShapeData, *outputShape, static_cast<size_t>(x_shape_size));
51+ } else {
52+ auto xShapeData = x_shape_tensor->GetData<int64_t>();
53+ return RandomUniformIntV2InferShapeImpl<int64_t>(xShapeData, *outputShape, static_cast<size_t>(x_shape_size));
54+ }
55+ 
56+ return ge::GRAPH_SUCCESS;
57+}
58+ 
59+IMPL_OP_INFERSHAPE(RandomUniformIntV2).InputsDataDependency({IN_SHAPE_IDX}).InferShape(InferShapeForRandomUniformIntV2);
60+} // namespace ops
Arandom/random_uniform_int_v2/op_kernel/arch35/random_uniform_int_v2.h+284-0
@@ -0,0 +1,284 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * 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+#pragma once
12+ 
13+#include "kernel_operator.h"
14+#include "op_kernel/platform_util.h"
15+#include "op_kernel/math_util.h"
16+ 
17+namespace RandomUniformIntV2 {
18+using namespace AscendC;
19+ 
20+template <typename T>
21+class RandomUniformIntV2Op {
22+public:
23+ __aicore__ inline RandomUniformIntV2Op(
24+ TPipe* pipe, const RandomUniformIntV2TilingData4RegBase* __restrict tilingData)
25+ : pipe_(pipe), tiling_(tilingData){};
26+ __aicore__ inline void Init(GM_ADDR y, GM_ADDR offset);
27+ __aicore__ inline void Process();
28+ 
29+private:
30+ __aicore__ inline void Skip(const uint64_t count);
31+ __aicore__ inline void DataTypeHandle(const uint32_t calCount);
32+ template <typename U, typename V>
33+ __aicore__ inline void UintToInt(LocalTensor<T>& yOutput, const uint32_t calCount);
34+ __aicore__ inline void CopyOut(int64_t yOffset, int64_t yCount);
35+ __aicore__ inline void offsetCopyOut(int64_t offsetValue);
36+ __aicore__ inline void InitKeyAndCounter();
37+ 
38+private:
39+ TPipe* pipe_;
40+ const RandomUniformIntV2TilingData4RegBase* tiling_;
41+ 
42+ constexpr static int8_t SAT_POS = 60;
43+ static constexpr uint16_t BUFFER_NUM = 2;
44+ static constexpr uint16_t ALG_KEY_SIZE = 2;
45+ static constexpr uint16_t ALG_COUNTER_SIZE = 4;
46+ static constexpr uint16_t RESULT_ELEMENT_CNT = 4;
47+ static constexpr uint16_t INT64_RATIO = 2;
48+ static constexpr uint64_t K_RESERVEED_PER_OUTPUT = 256;
49+ static constexpr uint32_t RIGHT_SHIFT = 32;
50+ 
51+ GlobalTensor<T> outputGm_;
52+ GlobalTensor<int64_t> offsetGm_;
53+ TBuf<QuePosition::VECCALC> philoxQueBuf_;
54+ TBuf<QuePosition::VECCALC> offsetBuf_;
55+ TQue<QuePosition::VECOUT, BUFFER_NUM> outQueY_;
56+ 
57+ int64_t curCoreProNum_ = 0;
58+ uint32_t key_[ALG_KEY_SIZE] = {0};
59+ uint32_t counter_[ALG_COUNTER_SIZE] = {0};
60+ 
61+ uint32_t blockIdx_;
62+ static constexpr MicroAPI::CastTrait castTraitTf = {
63+ MicroAPI::RegLayout::ZERO, MicroAPI::SatMode::NO_SAT, MicroAPI::MaskMergeMode::ZEROING};
64+};
65+ 
66+constexpr uint32_t THREAD_DIM = 512;
67+ 
68+template <typename T, typename UINT_T>
69+__simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_DIM) inline void UintToIntSimt(
70+ const uint32_t calCount, T low, UINT_T range, const __ubuf__ UINT_T* philox, __ubuf__ T* yOutput)
71+{
72+ for (uint32_t index = Simt::GetThreadIdx(); index < calCount; index = index + Simt::GetThreadNum()) {
73+ UINT_T randomRes = philox[index];
74+ UINT_T divRes = randomRes / range;
75+ UINT_T b = randomRes - divRes * range;
76+ UINT_T b_div_2 = b >> 1;
77+ yOutput[index] = low + static_cast<T>(b_div_2) + static_cast<T>(b - b_div_2);
78+ }
79+}
80+ 
81+template <typename T>
82+__aicore__ inline void RandomUniformIntV2Op<T>::Init(GM_ADDR y, GM_ADDR offset)
83+{
84+ blockIdx_ = GetBlockIdx();
85+ if (blockIdx_ > tiling_->blockNum) {
86+ return;
87+ }
88+ 
89+ if (blockIdx_ == tiling_->blockNum - 1) {
90+ curCoreProNum_ = tiling_->tailCoreProNum;
91+ } else {
92+ curCoreProNum_ = tiling_->normalCoreProNum;
93+ }
94+ 
95+ outputGm_.SetGlobalBuffer((__gm__ T*)y);
96+ offsetGm_.SetGlobalBuffer((__gm__ int64_t*)offset);
97+ pipe_->InitBuffer(outQueY_, BUFFER_NUM, tiling_->singleUbSize * sizeof(T));
98+ pipe_->InitBuffer(philoxQueBuf_, tiling_->singleUbSize * sizeof(uint32_t));
99+ pipe_->InitBuffer(offsetBuf_, Ops::Base::GetUbBlockSize());
100+}
101+ 
102+template <typename T>
103+__aicore__ inline void RandomUniformIntV2Op<T>::InitKeyAndCounter()
104+{
105+ key_[0] = static_cast<uint32_t>(tiling_->seed);
106+ key_[1] = static_cast<uint32_t>(tiling_->seed >> RIGHT_SHIFT);
107+ counter_[0] = 0;
108+ counter_[1] = 0;
109+ counter_[2] = static_cast<uint32_t>(tiling_->seed2);
110+ counter_[3] = static_cast<uint32_t>(tiling_->seed2 >> RIGHT_SHIFT);
111+}
112+ 
113+template <typename T>
114+__aicore__ inline void RandomUniformIntV2Op<T>::Process()
115+{
116+ if (blockIdx_ > tiling_->blockNum) {
117+ return;
118+ }
119+ 
120+ InitKeyAndCounter();
121+ auto offsetValue = offsetGm_.GetValue(0);
122+ if (offsetValue > 0) {
123+ Skip(offsetValue);
124+ }
125+ SyncAll();
126+ if (blockIdx_ == 0) {
127+ if (offsetValue < 0) {
128+ offsetValue = 0;
129+ }
130+ offsetValue = offsetValue + tiling_->outputSize * K_RESERVEED_PER_OUTPUT;
131+ offsetCopyOut(offsetValue);
132+ }
133+ uint16_t dtypeRatio = 1;
134+ if constexpr (AscendC::IsSameType<T, int64_t>::value) {
135+ dtypeRatio = INT64_RATIO;
136+ }
137+ auto blockOffSet = tiling_->normalCoreProNum * blockIdx_;
138+ auto resultElementCnt = RESULT_ELEMENT_CNT / dtypeRatio;
139+ auto groupCnt = Ops::Base::CeilDiv(blockOffSet, static_cast<int64_t>(resultElementCnt));
140+ 
141+ Skip(groupCnt);
142+ int64_t singleUbEleNum = tiling_->singleUbSize;
143+ int64_t ubRepeatimes = Ops::Base::CeilDiv(curCoreProNum_, singleUbEleNum);
144+ for (auto idx = 0; idx < ubRepeatimes; idx++) {
145+ int64_t curUbEleNum =
146+ idx == (ubRepeatimes - 1) ? curCoreProNum_ - (ubRepeatimes - 1) * singleUbEleNum : singleUbEleNum;
147+ int64_t philoxNumPro = curUbEleNum * dtypeRatio;
148+ int64_t philoxNumOffset = idx * singleUbEleNum;
149+ 
150+ LocalTensor<uint32_t> philoxRes = philoxQueBuf_.Get<uint32_t>();
151+ PhiloxRandom<10>(
152+ philoxRes, {key_[0], key_[1]}, {counter_[0], counter_[1], counter_[2], counter_[3]}, philoxNumPro);
153+ 
154+ DataTypeHandle(curUbEleNum);
155+ int64_t yOffset = blockOffSet + philoxNumOffset;
156+ CopyOut(yOffset, curUbEleNum);
157+ groupCnt = Ops::Base::CeilDiv(curUbEleNum, static_cast<int64_t>(resultElementCnt));
158+ Skip(groupCnt);
159+ }
160+}
161+ 
162+template <typename T>
163+__aicore__ inline void RandomUniformIntV2Op<T>::Skip(const uint64_t count)
164+{
165+ const uint32_t countLo = static_cast<uint32_t>(count);
166+ uint32_t countHi = static_cast<uint32_t>(count >> RIGHT_SHIFT);
167+ 
168+ counter_[0] += countLo;
169+ if (counter_[0] < countLo) {
170+ ++countHi;
171+ }
172+ counter_[1] += countHi;
173+ if (counter_[1] < countHi) {
174+ if (++counter_[2] == 0) {
175+ ++counter_[3];
176+ }
177+ }
178+}
179+ 
180+template <typename T>
181+__aicore__ inline void RandomUniformIntV2Op<T>::DataTypeHandle(const uint32_t calCount)
182+{
183+ LocalTensor<T> yOutput = outQueY_.AllocTensor<T>();
184+ if constexpr (AscendC::IsSameType<T, int32_t>::value) {
185+ UintToInt<int32_t, uint32_t>(yOutput, calCount);
186+ } else if constexpr (AscendC::IsSameType<T, int64_t>::value) {
187+ LocalTensor<uint64_t> philoxRes = philoxQueBuf_.Get<uint64_t>();
188+ AscendC::Simt::VF_CALL<UintToIntSimt<int64_t, uint64_t>>(
189+ AscendC::Simt::Dim3{THREAD_DIM}, calCount, tiling_->lo, tiling_->range,
190+ (__ubuf__ uint64_t*)philoxRes.GetPhyAddr(), (__ubuf__ int64_t*)yOutput.GetPhyAddr());
191+ }
192+ outQueY_.EnQue(yOutput);
193+}
194+ 
195+template <typename T>
196+__aicore__ inline void RandomUniformIntV2Op<T>::offsetCopyOut(int64_t offsetValue)
197+{
198+ LocalTensor<int64_t> offsetOutput = offsetBuf_.Get<int64_t>();
199+ offsetOutput.SetValue(0, offsetValue);
200+ DataCopyExtParams copyParams;
201+ copyParams.blockCount = 1;
202+ copyParams.blockLen = static_cast<uint32_t>(sizeof(int64_t));
203+ event_t SToMTE3Event = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::S_MTE3));
204+ SetFlag<HardEvent::S_MTE3>(SToMTE3Event);
205+ WaitFlag<HardEvent::S_MTE3>(SToMTE3Event);
206+ DataCopyPad(offsetGm_[0], offsetOutput, copyParams);
207+}
208+ 
209+template <typename T>
210+__aicore__ inline void RandomUniformIntV2Op<T>::CopyOut(int64_t yOffset, int64_t yCount)
211+{
212+ LocalTensor<T> yOutput = outQueY_.DeQue<T>();
213+ __ubuf__ T* ubPhilox = (__ubuf__ T*)yOutput.GetPhyAddr();
214+ DataCopyExtParams copyParams;
215+ copyParams.blockCount = 1;
216+ copyParams.blockLen = static_cast<uint32_t>(yCount * sizeof(T));
217+ DataCopyPad(outputGm_[yOffset], yOutput, copyParams);
218+ outQueY_.FreeTensor(yOutput);
219+}
220+ 
221+/*
222+ 将uint随机数转化生成[minVal, minVal + range)内整数
223+ 计算逻辑
224+ 1. 对Philox算法生成的随机数取模运算:b = philoxRandom % range
225+ 2. 计算第一步结果/2的结果暂存在b_div_2:b_div_2 = b >> 1
226+ 3. 获得最终输出:res = minVal + (int)b_div_2 + (int)(b-b_div_2)
227+*/
228+template <typename T>
229+template <typename U, typename V>
230+__aicore__ inline void RandomUniformIntV2Op<T>::UintToInt(LocalTensor<T>& yOutput, const uint32_t calCount)
231+{
232+ LocalTensor<V> philoxRes = philoxQueBuf_.Get<V>();
233+ __ubuf__ V* ubPhilox = (__ubuf__ V*)philoxRes.GetPhyAddr();
234+ __ubuf__ U* ubOut = (__ubuf__ U*)yOutput.GetPhyAddr();
235+ U minVal = tiling_->lo;
236+ V range = tiling_->range;
237+ uint32_t repeatTimes = Ops::Base::CeilDiv(calCount, static_cast<uint32_t>(Ops::Base::GetVRegSize() / sizeof(U)));
238+ 
239+ SetCtrlSpr<SAT_POS, SAT_POS>(0);
240+ __VEC_SCOPE__
241+ {
242+ MicroAPI::RegTensor<V> vReg0;
243+ MicroAPI::RegTensor<U> vReg1;
244+ MicroAPI::RegTensor<V> vReg2;
245+ MicroAPI::RegTensor<U> vReg3;
246+ MicroAPI::RegTensor<U> vReg4;
247+ MicroAPI::RegTensor<U> vReg5;
248+ MicroAPI::RegTensor<V> vReg6;
249+ MicroAPI::RegTensor<V> vReg7;
250+ MicroAPI::RegTensor<V> vReg8;
251+ MicroAPI::MaskReg mask;
252+ 
253+ uint32_t sReg1 = static_cast<uint32_t>(calCount);
254+ V sReg2 = range;
255+ U sReg3 = minVal;
256+ int16_t sReg4 = 1;
257+ 
258+ MicroAPI::MaskReg maskAll = MicroAPI::CreateMask<U, MicroAPI::MaskPattern::ALL>();
259+ MicroAPI::Duplicate<U, MicroAPI::MaskMergeMode::ZEROING>(vReg1, minVal, maskAll);
260+ MicroAPI::Duplicate<V, MicroAPI::MaskMergeMode::ZEROING>(vReg8, sReg2, maskAll);
261+ 
262+ int32_t offSet = static_cast<int32_t>(Ops::Base::GetVRegSize() / sizeof(T));
263+ for (uint16_t i = 0; i < static_cast<uint16_t>(repeatTimes); ++i) {
264+ mask = MicroAPI::UpdateMask<U>(sReg1);
265+ MicroAPI::DataCopy<V, MicroAPI::PostLiteral::POST_MODE_UPDATE, MicroAPI::LoadDist::DIST_NORM>(
266+ vReg0, ubPhilox, offSet);
267+ MicroAPI::Div<V, MicroAPI::MaskMergeMode::ZEROING>(vReg2, vReg0, vReg8, mask);
268+ MicroAPI::Mul<V, MicroAPI::MaskMergeMode::ZEROING>(vReg2, vReg2, vReg8, mask);
269+ MicroAPI::Sub<V, MicroAPI::MaskMergeMode::ZEROING>(vReg2, vReg0, vReg2, mask);
270+ MicroAPI::ShiftRights<V, int16_t>(vReg6, vReg2, sReg4, mask);
271+ MicroAPI::Sub<V, MicroAPI::MaskMergeMode::ZEROING>(vReg7, vReg2, vReg6, mask);
272+ vReg4 = (MicroAPI::RegTensor<U>&)vReg6;
273+ vReg5 = (MicroAPI::RegTensor<U>&)vReg7;
274+ MicroAPI::Add(vReg3, vReg1, vReg4, mask);
275+ MicroAPI::Add(vReg3, vReg3, vReg5, mask);
276+ 
277+ MicroAPI::DataCopy<U, MicroAPI::PostLiteral::POST_MODE_UPDATE, MicroAPI::StoreDist::DIST_NORM_B32>(
278+ ubOut, vReg3, offSet, mask);
279+ }
280+ }
281+ SetCtrlSpr<SAT_POS, SAT_POS>(1);
282+}
283+ 
284+} // namespace RandomUniformIntV2
Arandom/random_uniform_int_v2/op_kernel/arch35/random_uniform_int_v2_struct.h+44-0
@@ -0,0 +1,44 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * 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 random_uniform_int_v2_struct.h
13+ * \brief tiling base data
14+ */
15+ 
16+#ifndef Random_Uniform_Int_V2_STRUCT_H
17+#define Random_Uniform_Int_V2_STRUCT_H
18+ 
19+#include "ascendc/host_api/tiling/template_argument.h"
20+ 
21+class RandomUniformIntV2TilingData4RegBase {
22+public:
23+ int64_t blockNum = 0;
24+ int64_t normalCoreProNum = 0;
25+ int64_t tailCoreProNum = 0;
26+ int64_t singleUbSize = 0;
27+ int64_t seed = 0;
28+ int64_t seed2 = 0;
29+ int64_t outputSize = 0;
30+ uint64_t range = 0;
31+ int64_t lo = 0;
32+};
33+ 
34+#define RANDOM_UNIFORM_INT_V2_TPL 0
35+ 
36+ASCENDC_TPL_ARGS_DECL(RandomUniformIntV2,
37+ ASCENDC_TPL_UINT_DECL(opType, 1, ASCENDC_TPL_UI_LIST, RANDOM_UNIFORM_INT_V2_TPL)
38+);
39+ 
40+ASCENDC_TPL_SEL(
41+ ASCENDC_TPL_ARGS_SEL(
42+ ASCENDC_TPL_UINT_SEL(opType, ASCENDC_TPL_UI_LIST, RANDOM_UNIFORM_INT_V2_TPL)));
43+ 
44+#endif // Random_Uniform_Int_V2_STRUCT_H
Arandom/random_uniform_int_v2/op_kernel/random_uniform_int_v2.cpp+39-0
@@ -0,0 +1,39 @@
1+/**
2+* Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* 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 random_uniform_int_v2.cpp
13+* \brief
14+*/
15+ 
16+#include "arch35/random_uniform_int_v2_struct.h"
17+#include "arch35/random_uniform_int_v2.h"
18+using namespace AscendC;
19+using namespace RandomUniformIntV2;
20+ 
21+enum class RandomUniformIntV2TilingKey : uint32_t
22+{
23+ RANDOM_UNIFORM_INT = 0,
24+};
25+ 
26+template <uint32_t opType>
27+__global__ __aicore__ void random_uniform_int_v2(GM_ADDR shape, GM_ADDR min, GM_ADDR max, GM_ADDR inOffset, GM_ADDR y, GM_ADDR outOffset, GM_ADDR workspace, GM_ADDR tiling)
28+{
29+ REGISTER_TILING_DEFAULT(RandomUniformIntV2TilingData4RegBase);
30+ GET_TILING_DATA_WITH_STRUCT(RandomUniformIntV2TilingData4RegBase, tilingData, tiling);
31+ KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
32+ 
33+ AscendC::TPipe pipe;
34+ if constexpr(opType == static_cast<uint32_t>(RandomUniformIntV2TilingKey::RANDOM_UNIFORM_INT)) {
35+ RandomUniformIntV2::RandomUniformIntV2Op<DTYPE_Y> op(&pipe, &tilingData);
36+ op.Init(y, outOffset);
37+ op.Process();
38+ }
39+}
Arandom/random_uniform_int_v2/tests/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+foreach(SUB_DIR ${CURRENT_DIRS})
13+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14+ add_subdirectory(${SUB_DIR})
15+ endif()
16+endforeach()
Arandom/random_uniform_int_v2/tests/ut/CMakeLists.txt+16-0
@@ -0,0 +1,16 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
12+foreach(SUB_DIR ${CURRENT_DIRS})
13+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
14+ add_subdirectory(${SUB_DIR})
15+ endif()
16+endforeach()
Arandom/random_uniform_int_v2/tests/ut/op_host/CMakLists.txt+21-0
@@ -0,0 +1,21 @@
1+# ----------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------
10+ 
11+if(UT_TEST_ALL OR OP_HOST_UT)
12+ add_modules_ut_sources(UT_NAME ${OP_TILING_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
13+ add_modules_ut_sources(UT_NAME ${OP_INFERSHAPE_MODULE_NAME} MODE PRIVATE DIR ${CMAKE_CURRENT_SOURCE_DIR})
14+endif()
15+ 
16+file(GLOB CURRENT_DIRS RELATIVE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}/*)
17+foreach(SUB_DIR ${CURRENT_DIRS})
18+ if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/${SUB_DIR}/CMakeLists.txt")
19+ add_subdirectory(${SUB_DIR})
20+ endif()
21+endforeach()
Arandom/random_uniform_int_v2/tests/ut/op_host/arch35/test_random_uniform_int_v2_tiling.cpp+189-0
@@ -0,0 +1,189 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <iostream>
12+#include <gtest/gtest.h>
13+#include "tiling_context_faker.h"
14+#include "tiling_case_executor.h"
15+#include "../../../../op_host/arch35/random_uniform_int_v2_tiling_arch35.h"
16+ 
17+using namespace std;
18+using namespace ge;
19+ 
20+class RandomUniformIntV2Tiling : public testing::Test {
21+ protected:
22+ static void SetUpTestCase() {
23+ std::cout << "RandomUniformIntV2 SetUp" << std::endl;
24+ }
25+ 
26+ static void TearDownTestCase() {
27+ std::cout << "RandomUniformIntV2 TearDown" << std::endl;
28+ }
29+};
30+ 
31+TEST_F(RandomUniformIntV2Tiling, random_uniform_int_v2_tiling_950_int32_int32)
32+{
33+ optiling::RandomUniformIntV2CompileInfo compileInfo = {64, 196608};
34+ gert::StorageShape shape_shape = {{2}, {2}};
35+ gert::StorageShape min_shape = {{1}, {1}};
36+ gert::StorageShape max_shape = {{1}, {1}};
37+ gert::StorageShape offset_shape = {{1}, {1}};
38+ gert::StorageShape out_shape = {{32, 512}, {32, 512}};
39+ auto seed = Ops::Math::AnyValue::CreateFrom<int64_t>(10);
40+ auto seed2 = Ops::Math::AnyValue::CreateFrom<int64_t>(5);
41+ 
42+ vector<int32_t> shape_value = {32, 512};
43+ vector<int32_t> min_value = {2147483646};
44+ vector<int32_t> max_value = {2147483647};
45+ vector<int64_t> offset_value = {0};
46+ 
47+ gert::TilingContextPara tilingContextPara(
48+ "RandomUniformIntV2",
49+ {
50+ {shape_shape, ge::DT_INT32, ge::FORMAT_ND, true, shape_value.data()},
51+ {min_shape, ge::DT_INT32, ge::FORMAT_ND, true, min_value.data()},
52+ {max_shape, ge::DT_INT32, ge::FORMAT_ND, true, max_value.data()},
53+ {offset_shape, ge::DT_INT64, ge::FORMAT_ND, true, offset_value.data()}
54+ },
55+ {
56+ {out_shape, ge::DT_INT32, ge::FORMAT_ND},
57+ {offset_shape, ge::DT_INT64, ge::FORMAT_ND}
58+ },
59+ {
60+ gert::TilingContextPara::OpAttr("seed", seed),
61+ gert::TilingContextPara::OpAttr("seed2", seed2)
62+ },
63+ &compileInfo);
64+ uint64_t expectTilingKey = 0;
65+ string expectTilingData =
66+ "64 256 256 12288 10 5 16384 1 2147483646 ";
67+ std::vector<size_t> expectWorkspaces = {0};
68+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
69+}
70+ 
71+TEST_F(RandomUniformIntV2Tiling, random_uniform_int_v2_tiling_950_int32_int64)
72+{
73+ optiling::RandomUniformIntV2CompileInfo compileInfo = {64, 196608};
74+ gert::StorageShape shape_shape = {{2}, {2}};
75+ gert::StorageShape min_shape = {{1}, {1}};
76+ gert::StorageShape max_shape = {{1}, {1}};
77+ gert::StorageShape offset_shape = {{1}, {1}};
78+ gert::StorageShape out_shape = {{32, 512}, {32, 512}};
79+ auto seed = Ops::Math::AnyValue::CreateFrom<int64_t>(10);
80+ auto seed2 = Ops::Math::AnyValue::CreateFrom<int64_t>(5);
81+ 
82+ vector<int32_t> shape_value = {32, 512};
83+ vector<int64_t> min_value = {2147483648};
84+ vector<int64_t> max_value = {2147483649};
85+ vector<int64_t> offset_value = {0};
86+ 
87+ gert::TilingContextPara tilingContextPara(
88+ "RandomUniformIntV2",
89+ {
90+ {shape_shape, ge::DT_INT32, ge::FORMAT_ND, true, shape_value.data()},
91+ {min_shape, ge::DT_INT64, ge::FORMAT_ND, true, min_value.data()},
92+ {max_shape, ge::DT_INT64, ge::FORMAT_ND, true, max_value.data()},
93+ {offset_shape, ge::DT_INT64, ge::FORMAT_ND, true, offset_value.data()}
94+ },
95+ {
96+ {out_shape, ge::DT_INT64, ge::FORMAT_ND},
97+ {offset_shape, ge::DT_INT64, ge::FORMAT_ND}
98+ },
99+ {
100+ gert::TilingContextPara::OpAttr("seed", seed),
101+ gert::TilingContextPara::OpAttr("seed2", seed2)
102+ },
103+ &compileInfo);
104+ uint64_t expectTilingKey = 0;
105+ string expectTilingData =
106+ "64 256 256 6144 10 5 16384 1 2147483648 ";
107+ std::vector<size_t> expectWorkspaces = {0};
108+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
109+}
110+ 
111+TEST_F(RandomUniformIntV2Tiling, random_uniform_int_v2_tiling_950_int64_int32)
112+{
113+ optiling::RandomUniformIntV2CompileInfo compileInfo = {64, 196608};
114+ gert::StorageShape shape_shape = {{2}, {2}};
115+ gert::StorageShape min_shape = {{1}, {1}};
116+ gert::StorageShape max_shape = {{1}, {1}};
117+ gert::StorageShape offset_shape = {{1}, {1}};
118+ gert::StorageShape out_shape = {{32, 512}, {32, 512}};
119+ auto seed = Ops::Math::AnyValue::CreateFrom<int64_t>(10);
120+ auto seed2 = Ops::Math::AnyValue::CreateFrom<int64_t>(5);
121+ 
122+ vector<int64_t> shape_value = {32, 512};
123+ vector<int32_t> min_value = {2};
124+ vector<int32_t> max_value = {5};
125+ vector<int64_t> offset_value = {0};
126+ 
127+ gert::TilingContextPara tilingContextPara(
128+ "RandomUniformIntV2",
129+ {
130+ {shape_shape, ge::DT_INT64, ge::FORMAT_ND, true, shape_value.data()},
131+ {min_shape, ge::DT_INT32, ge::FORMAT_ND, true, min_value.data()},
132+ {max_shape, ge::DT_INT32, ge::FORMAT_ND, true, max_value.data()},
133+ {offset_shape, ge::DT_INT64, ge::FORMAT_ND, true, offset_value.data()}
134+ },
135+ {
136+ {out_shape, ge::DT_INT32, ge::FORMAT_ND},
137+ {offset_shape, ge::DT_INT64, ge::FORMAT_ND}
138+ },
139+ {
140+ gert::TilingContextPara::OpAttr("seed", seed),
141+ gert::TilingContextPara::OpAttr("seed2", seed2)
142+ },
143+ &compileInfo);
144+ uint64_t expectTilingKey = 0;
145+ string expectTilingData =
146+ "64 256 256 12288 10 5 16384 3 2 ";
147+ std::vector<size_t> expectWorkspaces = {0};
148+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
149+}
150+ 
151+TEST_F(RandomUniformIntV2Tiling, random_uniform_int_v2_tiling_950_int64_int64)
152+{
153+ optiling::RandomUniformIntV2CompileInfo compileInfo = {64, 196608};
154+ gert::StorageShape shape_shape = {{2}, {2}};
155+ gert::StorageShape min_shape = {{1}, {1}};
156+ gert::StorageShape max_shape = {{1}, {1}};
157+ gert::StorageShape offset_shape = {{1}, {1}};
158+ gert::StorageShape out_shape = {{32, 512}, {32, 512}};
159+ auto seed = Ops::Math::AnyValue::CreateFrom<int64_t>(10);
160+ auto seed2 = Ops::Math::AnyValue::CreateFrom<int64_t>(5);
161+ 
162+ vector<int64_t> shape_value = {32, 512};
163+ vector<int64_t> min_value = {2};
164+ vector<int64_t> max_value = {5};
165+ vector<int64_t> offset_value = {0};
166+ 
167+ gert::TilingContextPara tilingContextPara(
168+ "RandomUniformIntV2",
169+ {
170+ {shape_shape, ge::DT_INT64, ge::FORMAT_ND, true, shape_value.data()},
171+ {min_shape, ge::DT_INT64, ge::FORMAT_ND, true, min_value.data()},
172+ {max_shape, ge::DT_INT64, ge::FORMAT_ND, true, max_value.data()},
173+ {offset_shape, ge::DT_INT64, ge::FORMAT_ND, true, offset_value.data()}
174+ },
175+ {
176+ {out_shape, ge::DT_INT64, ge::FORMAT_ND},
177+ {offset_shape, ge::DT_INT64, ge::FORMAT_ND}
178+ },
179+ {
180+ gert::TilingContextPara::OpAttr("seed", seed),
181+ gert::TilingContextPara::OpAttr("seed2", seed2)
182+ },
183+ &compileInfo);
184+ uint64_t expectTilingKey = 0;
185+ string expectTilingData =
186+ "64 256 256 6144 10 5 16384 3 2 ";
187+ std::vector<size_t> expectWorkspaces = {0};
188+ ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectTilingData, expectWorkspaces);
189+}
Arandom/random_uniform_int_v2/tests/ut/op_host/test_random_uniform_int_v2_infershape.cpp+139-0
@@ -0,0 +1,139 @@
1+/**
2+ * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include <gtest/gtest.h>
12+#include <iostream>
13+#include "infershape_context_faker.h"
14+#include "base/registry/op_impl_space_registry_v2.h"
15+ 
16+class random_uniform_int_v2 : public testing::Test {
17+protected:
18+ static void SetUpTestCase()
19+ {
20+ std::cout << "random_uniform_int_v2 SetUp" << std::endl;
21+ }
22+ 
23+ static void TearDownTestCase()
24+ {
25+ std::cout << "random_uniform_int_v2 TearDown" << std::endl;
26+ }
27+};
28+ 
29+static std::vector<int64_t> ToVector(const gert::Shape& shape)
30+{
31+ size_t shapeSize = shape.GetDimNum();
32+ std::vector<int64_t> shapeVec(shapeSize, 0);
33+ for (size_t i = 0; i < shapeSize; i++) {
34+ shapeVec[i] = shape.GetDim(i);
35+ }
36+ return shapeVec;
37+}
38+ 
39+template <typename T>
40+gert::Tensor* ConstructInputConstTensor(
41+ std::unique_ptr<uint8_t[]>& input_tensor_holder, const std::vector<T>& const_value, ge::DataType const_dtype)
42+{
43+ size_t shape_size = const_value.size();
44+ std::cout << " shape_size:" << shape_size << std::endl;
45+ auto input_tensor = reinterpret_cast<gert::Tensor*>(input_tensor_holder.get());
46+ gert::Tensor tensor(
47+ {{shape_size}, {shape_size}}, // shape
48+ {ge::FORMAT_ND, ge::FORMAT_ND, {}}, // format
49+ gert::kFollowing, // placement
50+ const_dtype, // dt
51+ nullptr);
52+ std::memcpy(input_tensor, &tensor, sizeof(gert::Tensor));
53+ auto tensor_data = reinterpret_cast<T*>(input_tensor + 1);
54+ for (size_t i = 0; i < shape_size; i++) {
55+ tensor_data[i] = const_value[i];
56+ std::cout << " const_value:" << const_value[i] << std::endl;
57+ }
58+ 
59+ input_tensor->SetData(gert::TensorData(tensor_data, nullptr));
60+ return input_tensor;
61+}
62+ 
63+static void ExeTestCase(
64+ std::vector<std::vector<int64_t>> expectResults,
65+ const std::vector<gert::StorageShape>& inputShapes, // 存储所有输入StorageShape参数
66+ const std::vector<ge::DataType>& dtypes, // 存储所有DataType参数
67+ const std::vector<int32_t> input1Values, const std::vector<int64_t> input2Values, const std::vector<int64_t> input3Values, const std::vector<int64_t> input4Values, std::vector<gert::StorageShape *>& outStorageShape,
68+ ge::graphStatus testCaseResult = ge::GRAPH_SUCCESS)
69+{
70+ // 从vector中取出对应参数(保持原顺序)
71+ const auto& shapeStorageShape = inputShapes[0];
72+ const auto& minStorageShape = inputShapes[1];
73+ const auto& maxStorageShape = inputShapes[2];
74+ const auto& offsetStorageShape = inputShapes[3];
75+ 
76+ ge::DataType input1Dtype = dtypes[0];
77+ ge::DataType input2Dtype = dtypes[1];
78+ ge::DataType input3Dtype = dtypes[2];
79+ 
80+ auto shape_input_tensor = std::unique_ptr<uint8_t[]>(new uint8_t[sizeof(gert::Tensor) + sizeof(int32_t) * input1Values.size()]);
81+ auto shape_tensor = ConstructInputConstTensor<int32_t>(shape_input_tensor, input1Values, ge::DT_INT32);
82+ auto min_input_tensor = std::unique_ptr<uint8_t[]>(new uint8_t[sizeof(gert::Tensor) + sizeof(int64_t) * input2Values.size()]);
83+ auto min_tensor = ConstructInputConstTensor<int64_t>(min_input_tensor, input2Values, ge::DT_INT64);
84+ auto max_input_tensor = std::unique_ptr<uint8_t[]>(new uint8_t[sizeof(gert::Tensor) + sizeof(int64_t) * input3Values.size()]);
85+ auto max_tensor = ConstructInputConstTensor<int64_t>(max_input_tensor, input3Values, ge::DT_INT64);
86+ auto offset_input_tensor = std::unique_ptr<uint8_t[]>(new uint8_t[sizeof(gert::Tensor) + sizeof(int64_t)]);
87+ auto offset_tensor = ConstructInputConstTensor<int64_t>(offset_input_tensor, input4Values, ge::DT_INT64);
88+ 
89+ /* make infershape context */
90+ std::vector<gert::Tensor*> inputTensors = {shape_tensor, min_tensor, max_tensor, offset_tensor};
91+ 
92+ std::vector<gert::StorageShape*> outputShapes = outStorageShape;
93+ auto contextHolder = gert::InferShapeContextFaker()
94+ .SetOpType("RandomUniformIntV2")
95+ .NodeIoNum(4, 2)
96+ .NodeInputTd(0, input1Dtype, ge::FORMAT_ND, ge::FORMAT_ND)
97+ .NodeInputTd(1, input2Dtype, ge::FORMAT_ND, ge::FORMAT_ND)
98+ .NodeInputTd(2, input2Dtype, ge::FORMAT_ND, ge::FORMAT_ND)
99+ .NodeInputTd(3, input3Dtype, ge::FORMAT_ND, ge::FORMAT_ND)
100+ .NodeOutputTd(0, input2Dtype, ge::FORMAT_ND, ge::FORMAT_ND)
101+ .NodeOutputTd(1, input3Dtype, ge::FORMAT_ND, ge::FORMAT_ND)
102+ .InputTensors(inputTensors)
103+ .OutputShapes(outputShapes)
104+ .Build();
105+ 
106+ /* get infershape func */
107+ auto spaceRegistry = gert::DefaultOpImplSpaceRegistryV2::GetInstance().GetSpaceRegistry();
108+ auto inferShapeFunc = spaceRegistry->GetOpImpl("RandomUniformIntV2")->infer_shape;
109+ ASSERT_NE(inferShapeFunc, nullptr);
110+ 
111+ /* do infershape */
112+ EXPECT_EQ(inferShapeFunc(contextHolder.GetContext()), testCaseResult);
113+ for (size_t i = 0; i < expectResults.size(); i++) {
114+ EXPECT_EQ(ToVector(*contextHolder.GetContext()->GetOutputShape(i)), expectResults[i]);
115+ }
116+}
117+ 
118+TEST_F(random_uniform_int_v2, random_uniform_int_v2_infershape_case_0)
119+{
120+ // 用vector存储同类型参数(顺序与原参数列表一致)
121+ std::vector<gert::StorageShape> inputShapes = {
122+ {{32, 512}, {32, 512}},
123+ {{1}, {1}},
124+ {{1}, {1}},
125+ {{1}, {1}},
126+ };
127+ std::vector<ge::DataType> dtypes = {ge::DT_INT32, ge::DT_INT64, ge::DT_INT64};
128+ 
129+ std::vector<int32_t> input1Values = {32, 512};
130+ std::vector<int64_t> input2Values = {1};
131+ std::vector<int64_t> input3Values = {10};
132+ std::vector<int64_t> input4Values = {0};
133+ std::vector<std::vector<int64_t>> expectResult = {{32, 512}, {1}};
134+ std::vector<gert::StorageShape*> outStorageShape = {};
135+ 
136+ // 简化后的函数调用
137+ ExeTestCase(
138+ expectResult, inputShapes, dtypes, input1Values, input2Values, input3Values, input4Values, outStorageShape, ge::GRAPH_SUCCESS);
139+}
Mscripts/ci/ascend950/ops_math_operator_list.yaml+1-0
@@ -163,6 +163,7 @@ operator_group_5:
163 - drop_out_v3163 - drop_out_v3
164 - random_standard_normal_v2164 - random_standard_normal_v2
165 - random_uniform_v2165 - random_uniform_v2
166+ - random_uniform_int_v2
166 - stateless_bernoulli167 - stateless_bernoulli
167 - stateless_drop_out_gen_mask168 - stateless_drop_out_gen_mask
168 - stateless_random_normal_v2169 - stateless_random_normal_v2