已合并
support RandomUniformIntV2 #957
xuejinghui创建于 1月27日
support RandomUniformIntV2 #957
已合并
从已删除 :master合入到cann/ops-mathmaster
共 23 个文件变更+2244-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> |
| @@ -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}) | ||
| @@ -0,0 +1,102 @@ | |||
| 1 | +# RandomUniformIntV2 | ||
C | |||
| 2 | + | ||
| 3 | +## 产品支持情况 | ||
| 4 | + | ||
| 5 | +| 产品 | 是否支持 | | ||
| 6 | +| :----------------------------------------------------------------- | :------: | | ||
| 7 | +| Ascend 950PR/Ascend 950DT | √ | | ||
| 8 | +| Atlas A3 训练系列产品/Atlas A3 推理系列产品 | × | | ||
| 9 | +| Atlas A2 训练系列产品/Atlas A2 推理系列产品 | × | | ||
C xuejinghui 1月30日 评论: 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算子。 | | ||
| @@ -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 | + | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + | ||
| 34 | + | ||
| 35 | + | ||
| 36 | +using namespace ge; | ||
| 37 | +using std::map; | ||
| 38 | +using std::string; | ||
| 39 | +using std::vector; | ||
| 40 | + | ||
| 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 | + | ||
| 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 | + | ||
| 86 | + add1.set_attr_##attrName(attrValue); | ||
| 87 | + | ||
| 88 | + | ||
| 89 | + TensorDesc outputName##outputIndex##_desc = \ | ||
| 90 | + TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \ | ||
| 91 | + add1.update_output_desc_##outputName(outputName##outputIndex##_desc); | ||
| 92 | + | ||
| 93 | + | ||
| 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 | + | ||
| 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 | +} | ||
| @@ -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() | ||
| @@ -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 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 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 | ||
| @@ -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 | ||
| @@ -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 | + | ||
| 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 | ||
| @@ -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 | + | ||
| 16 | + | ||
| 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 | ||
| @@ -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 | + | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 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 | ||
| @@ -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 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 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 | + | ||
| 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 | + | ||
| @@ -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 | + | ||
| 17 | + | ||
| 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 | +} | ||
| @@ -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() | ||
| @@ -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() | ||
| @@ -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() | ||
| @@ -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 | + | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 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 | +} | ||
| @@ -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 | + | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 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 | +} | ||
| @@ -163,6 +163,7 @@ operator_group_5: | |||
| 163 | - drop_out_v3 | 163 | - drop_out_v3 |
| 164 | - random_standard_normal_v2 | 164 | - random_standard_normal_v2 |
| 165 | - random_uniform_v2 | 165 | - random_uniform_v2 |
| 166 | + - random_uniform_int_v2 | ||
| 166 | - stateless_bernoulli | 167 | - stateless_bernoulli |
| 167 | - stateless_drop_out_gen_mask | 168 | - stateless_drop_out_gen_mask |
| 168 | - stateless_random_normal_v2 | 169 | - stateless_random_normal_v2 |


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