已合并
besseli1e算子开发 #3780
kangjiaming创建于 7月3日
besseli1e算子开发 #3780
已合并
共 19 个文件变更+1175-2
| @@ -356,6 +356,16 @@ | |||
| 356 | <td>AI Core</td> | 356 | <td>AI Core</td> |
| 357 | <td>计算指数缩放的第一类零阶修正贝塞尔函数，公式y = exp(-|x|) * I0(x)，仅支持图模式调用。</td> | 357 | <td>计算指数缩放的第一类零阶修正贝塞尔函数，公式y = exp(-|x|) * I0(x)，仅支持图模式调用。</td> |
| 358 | </tr> | 358 | </tr> |
| 359 | + <tr> | ||
| 360 | + <td>math</td> | ||
| 361 | + <td><a href="../../math/bessel_i1e/README.md">bessel_i1e</a></td> | ||
| 362 | + <td>√</td> | ||
| 363 | + <td>√</td> | ||
| 364 | + <td>×</td> | ||
| 365 | + <td>√</td> | ||
| 366 | + <td>AI Core</td> | ||
| 367 | + <td>计算指数缩放修正贝塞尔函数（第一类，阶数1）。</td> | ||
| 368 | + </tr> | ||
| 359 | <tr> | 369 | <tr> |
| 360 | <td>math</td> | 370 | <td>math</td> |
| 361 | <td><a href="../../math/bias_add/README.md">bias_add</a></td> | 371 | <td><a href="../../math/bias_add/README.md">bias_add</a></td> |
| @@ -0,0 +1,14 @@ | |||
| 1 | +#-- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- | ||
| 2 | +#Copyright(c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | +#This program is free software, you can redistribute it and / or modify it under the terms and conditions of | ||
| 4 | +#CANN Open Software License Agreement Version 2.0(the "License"). | ||
| 5 | +#Please refer to the License for details.You may not use this file except in compliance with the License. | ||
| 6 | +#THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | +#INCLUDING BUT NOT LIMITED TO NON - INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | +#See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | +#-- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- -- | ||
| 10 | +#Generated By CANNBot | ||
| 11 | + | ||
| 12 | +set(SUPPORT_COMPUTE_UNIT "ascend910b" "ascend910_93" "ascend950") | ||
| 13 | +set(SUPPORT_TILING_DIR "arch32" "arch32" "arch35") | ||
| 14 | +add_all_modules_sources(OPTYPE bessel_i1e ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE) | ||
| @@ -0,0 +1,72 @@ | |||
| 1 | +<!-- | ||
| 2 | +Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | +--> | ||
| 4 | + | ||
| 5 | +<!-- Generated By CANNBot --> | ||
| 6 | + | ||
| 7 | +# BesselI1e | ||
| 8 | + | ||
| 9 | +## 产品支持情况 | ||
| 10 | + | ||
| 11 | +| 产品 | 是否支持 | | ||
| 12 | +|:------------------------------------------------| :------: | | ||
| 13 | +| <term>Ascend 950PR/Ascend 950DT</term> | √ | | ||
| 14 | +| <term>Atlas A3 训练系列产品/Atlas A3 推理系列产品</term> | √ | | ||
| 15 | +| <term>Atlas A2 训练系列产品/Atlas A2 推理系列产品</term> | √ | | ||
| 16 | +| <term>Atlas 200I/500 A2 推理产品</term> | × | | ||
| 17 | +| <term>Atlas 推理系列产品</term> | × | | ||
| 18 | +| <term>Atlas 训练系列产品</term> | × | | ||
| 19 | + | ||
| 20 | +## 功能说明 | ||
| 21 | + | ||
| 22 | +- 算子功能：计算指数缩放的第一阶修正贝塞尔函数。 | ||
| 23 | + | ||
| 24 | +- 计算公式： | ||
| 25 | + | ||
| 26 | +$$out_i = e^{-|input_i|} \cdot I_1(input_i)$$ | ||
| 27 | + | ||
| 28 | +其中 $I_1(x)$ 是第一阶修正贝塞尔函数。指数缩放因子 $e^{-|x|}$ 使得函数在大参数值时数值稳定。 | ||
| 29 | + | ||
| 30 | +## 参数说明 | ||
| 31 | + | ||
| 32 | +<table style="undefined;table-layout: fixed; width: 980px"><colgroup> | ||
| 33 | + <col style="width: 100px"> | ||
| 34 | + <col style="width: 150px"> | ||
| 35 | + <col style="width: 280px"> | ||
| 36 | + <col style="width: 330px"> | ||
| 37 | + <col style="width: 120px"> | ||
| 38 | + </colgroup> | ||
| 39 | + <thead> | ||
| 40 | + <tr> | ||
| 41 | + <th>参数名</th> | ||
| 42 | + <th>输入/输出/属性</th> | ||
| 43 | + <th>描述</th> | ||
| 44 | + <th>数据类型</th> | ||
| 45 | + <th>数据格式</th> | ||
| 46 | + </tr></thead> | ||
| 47 | + <tbody> | ||
| 48 | + <tr> | ||
| 49 | + <td>x</td> | ||
| 50 | + <td>输入</td> | ||
| 51 | + <td>待进行BesselI1e计算的入参。</td> | ||
| 52 | + <td>FLOAT、FLOAT16、BF16</td> | ||
| 53 | + <td>ND</td> | ||
| 54 | + </tr> | ||
| 55 | + <tr> | ||
| 56 | + <td>y</td> | ||
| 57 | + <td>输出</td> | ||
| 58 | + <td>BesselI1e计算的出参。</td> | ||
| 59 | + <td>FLOAT、FLOAT16、BF16</td> | ||
| 60 | + <td>ND</td> | ||
| 61 | + </tr> | ||
| 62 | + </tbody></table> | ||
| 63 | + | ||
| 64 | +## 约束说明 | ||
| 65 | + | ||
| 66 | +无 | ||
| 67 | + | ||
| 68 | +## 调用说明 | ||
| 69 | + | ||
| 70 | +| 调用方式 | 调用样例 | 说明 | | ||
| 71 | +|:---------|:---------|:-----| | ||
| 72 | +| 图模式调用 | [test_geir_bessel_i1e](./examples/test_geir_bessel_i1e.cpp) | 通过[算子IR](./op_graph/bessel_i1e_proto.h)构图方式调用BesselI1e算子。 | | ||
| @@ -0,0 +1,197 @@ | |||||||||
| 1 | +/** | ||||||||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||||||||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||||||||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||||||||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||||||||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||||||||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||||||||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||||||||
A | |||||||||
| 9 | + * | ||||||||
| 10 | + * Generated By CANNBot | ||||||||
| 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 | + | ||||||||
| 38 | +static std::string GetTime() | ||||||||
| 39 | +{ | ||||||||
| 40 | + time_t timep; | ||||||||
| 41 | + time(&timep); | ||||||||
| 42 | + char tmp[64]; | ||||||||
| 43 | + strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep)); | ||||||||
| 44 | + return tmp; | ||||||||
| 45 | +} | ||||||||
| 46 | + | ||||||||
| 47 | +static uint32_t GetDataTypeSize(DataType dt) | ||||||||
| 48 | +{ | ||||||||
| 49 | + switch (dt) { | ||||||||
| 50 | + case DT_FLOAT: | ||||||||
| 51 | + return 4; | ||||||||
| 52 | + case DT_FLOAT16: | ||||||||
| 53 | + return 2; | ||||||||
| 54 | + default: | ||||||||
| 55 | + return 4; | ||||||||
| 56 | + } | ||||||||
| 57 | +} | ||||||||
| 58 | + | ||||||||
| 59 | +static int32_t GenTestData(const std::vector<int64_t>& shapes, Tensor& tensor, TensorDesc& desc, DataType dtype) | ||||||||
| 60 | +{ | ||||||||
| 61 | + desc.SetRealDimCnt(shapes.size()); | ||||||||
| 62 | + size_t size = 1; | ||||||||
| 63 | + for (auto d : shapes) | ||||||||
| 64 | + size *= d; | ||||||||
| 65 | + | ||||||||
| 66 | + size_t data_len = size * GetDataTypeSize(dtype); | ||||||||
| 67 | + float* pData = new (std::nothrow) float[size]; | ||||||||
| 68 | + if (pData == nullptr) | ||||||||
| 69 | + return FAILED; | ||||||||
| 70 | + for (size_t i = 0; i < size; ++i) { | ||||||||
| 71 | + pData[i] = static_cast<float>(i % 10) * 0.5f; | ||||||||
 🟡 Medium Priority
虽然对于该示例中 shape 建议：在 new 之后增加空指针检查： 改动建议
  不准确？ | |||||||||
| 72 | + } | ||||||||
| 73 | + tensor = Tensor(desc, reinterpret_cast<uint8_t*>(pData), data_len); | ||||||||
| 74 | + delete[] pData; | ||||||||
| 75 | + return SUCCESS; | ||||||||
| 76 | +} | ||||||||
| 77 | + | ||||||||
| 78 | +static int32_t WriteDataToFile(const std::string& bin_file, uint64_t data_size, uint8_t* data) | ||||||||
| 79 | +{ | ||||||||
| 80 | + FILE* fp = fopen(bin_file.c_str(), "wb"); | ||||||||
| 81 | + if (fp == nullptr) | ||||||||
| 82 | + return FAILED; | ||||||||
| 83 | + fwrite(data, sizeof(uint8_t), data_size, fp); | ||||||||
| 84 | + fclose(fp); | ||||||||
| 85 | + return SUCCESS; | ||||||||
| 86 | +} | ||||||||
| 87 | + | ||||||||
| 88 | +int CreateOpInGraph(DataType inDtype, std::vector<ge::Tensor>& input, std::vector<Operator>& inputs, | ||||||||
| 89 | + std::vector<Operator>& outputs, Graph& graph) | ||||||||
| 90 | +{ | ||||||||
| 91 | + Status ret = SUCCESS; | ||||||||
| 92 | + auto bessel_op = op::BesselI1e("bessel_i1e_0"); | ||||||||
| 93 | + | ||||||||
| 94 | + std::vector<int64_t> xShape = {32, 4, 4, 4}; | ||||||||
| 95 | + auto placeholder0 = op::Data("placeholder0").set_attr_index(0); | ||||||||
| 96 | + TensorDesc placeholder0_desc = TensorDesc(ge::Shape(xShape), FORMAT_ND, inDtype); | ||||||||
| 97 | + placeholder0_desc.SetPlacement(ge::kPlacementHost); | ||||||||
| 98 | + placeholder0_desc.SetFormat(FORMAT_ND); | ||||||||
| 99 | + Tensor tensor_placeholder0; | ||||||||
| 100 | + ret = GenTestData(xShape, tensor_placeholder0, placeholder0_desc, inDtype); | ||||||||
| 101 | + if (ret != SUCCESS) { | ||||||||
| 102 | + printf("%s - ERROR: Generate input data failed\n", GetTime().c_str()); | ||||||||
| 103 | + return FAILED; | ||||||||
| 104 | + } | ||||||||
| 105 | + placeholder0.update_input_desc_x(placeholder0_desc); | ||||||||
| 106 | + input.push_back(tensor_placeholder0); | ||||||||
| 107 | + graph.AddOp(placeholder0); | ||||||||
| 108 | + bessel_op.set_input_x(placeholder0); | ||||||||
| 109 | + | ||||||||
| 110 | + TensorDesc y_desc = TensorDesc(ge::Shape(xShape), FORMAT_ND, inDtype); | ||||||||
| 111 | + bessel_op.update_output_desc_y(y_desc); | ||||||||
| 112 | + | ||||||||
| 113 | + inputs.push_back(placeholder0); | ||||||||
| 114 | + outputs.push_back(bessel_op); | ||||||||
| 115 | + return SUCCESS; | ||||||||
| 116 | +} | ||||||||
| 117 | + | ||||||||
| 118 | +int main(int argc, char* argv[]) | ||||||||
| 119 | +{ | ||||||||
| 120 | + const char* graph_name = "bessel_i1e_geir_test"; | ||||||||
| 121 | + Graph graph(graph_name); | ||||||||
| 122 | + std::vector<ge::Tensor> input; | ||||||||
| 123 | + | ||||||||
| 124 | + printf("%s - INFO: Start to initialize GE\n", GetTime().c_str()); | ||||||||
| 125 | + std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}}; | ||||||||
| 126 | + Status ret = ge::GEInitialize(global_options); | ||||||||
| 127 | + if (ret != SUCCESS) { | ||||||||
| 128 | + printf("%s - ERROR: GE initialize failed\n", GetTime().c_str()); | ||||||||
| 129 | + return FAILED; | ||||||||
| 130 | + } | ||||||||
| 131 | + | ||||||||
| 132 | + std::vector<Operator> inputs{}; | ||||||||
| 133 | + std::vector<Operator> outputs{}; | ||||||||
| 134 | + DataType inDtype = DT_FLOAT; | ||||||||
| 135 | + | ||||||||
| 136 | + ret = CreateOpInGraph(inDtype, input, inputs, outputs, graph); | ||||||||
| 137 | + if (ret != SUCCESS) { | ||||||||
| 138 | + printf("%s - ERROR: Create op in graph failed\n", GetTime().c_str()); | ||||||||
| 139 | + ge::GEFinalize(); | ||||||||
| 140 | + return FAILED; | ||||||||
| 141 | + } | ||||||||
| 142 | + | ||||||||
| 143 | + if (!inputs.empty() && !outputs.empty()) { | ||||||||
| 144 | + graph.SetInputs(inputs).SetOutputs(outputs); | ||||||||
| 145 | + } | ||||||||
| 146 | + | ||||||||
| 147 | + std::map<AscendString, AscendString> build_options = {}; | ||||||||
| 148 | + printf("%s - INFO: Creating session\n", GetTime().c_str()); | ||||||||
| 149 | + ge::Session* session = new Session(build_options); | ||||||||
| 150 | + if (session == nullptr) { | ||||||||
| 151 | + printf("%s - ERROR: Create session failed\n", GetTime().c_str()); | ||||||||
| 152 | + ge::GEFinalize(); | ||||||||
| 153 | + return FAILED; | ||||||||
| 154 | + } | ||||||||
| 155 | + | ||||||||
| 156 | + std::map<AscendString, AscendString> graph_options = {}; | ||||||||
| 157 | + uint32_t graph_id = 0; | ||||||||
| 158 | + ret = session->AddGraph(graph_id, graph, graph_options); | ||||||||
| 159 | + if (ret != SUCCESS) { | ||||||||
| 160 | + printf("%s - ERROR: Add graph failed\n", GetTime().c_str()); | ||||||||
| 161 | + delete session; | ||||||||
| 162 | + ge::GEFinalize(); | ||||||||
| 163 | + return FAILED; | ||||||||
| 164 | + } | ||||||||
| 165 | + | ||||||||
| 166 | + printf("%s - INFO: Running graph\n", GetTime().c_str()); | ||||||||
| 167 | + std::vector<ge::Tensor> output; | ||||||||
| 168 | + ret = session->RunGraph(graph_id, input, output); | ||||||||
| 169 | + if (ret != SUCCESS) { | ||||||||
| 170 | + printf("%s - ERROR: Run graph failed\n", GetTime().c_str()); | ||||||||
| 171 | + delete session; | ||||||||
| 172 | + ge::GEFinalize(); | ||||||||
| 173 | + return FAILED; | ||||||||
| 174 | + } | ||||||||
| 175 | + printf("%s - INFO: Run graph success\n", GetTime().c_str()); | ||||||||
| 176 | + | ||||||||
| 177 | + for (size_t i = 0; i < output.size(); i++) { | ||||||||
| 178 | + std::string output_file = "./bessel_i1e_geir_output_" + std::to_string(i) + ".bin"; | ||||||||
| 179 | + uint8_t* output_data = output[i].GetData(); | ||||||||
| 180 | + int64_t output_shape_size = output[i].GetTensorDesc().GetShape().GetShapeSize(); | ||||||||
| 181 | + uint64_t data_size = static_cast<uint64_t>(output_shape_size) * | ||||||||
| 182 | + GetDataTypeSize(output[i].GetTensorDesc().GetDataType()); | ||||||||
| 183 | + WriteDataToFile(output_file, data_size, output_data); | ||||||||
| 184 | + printf("%s - INFO: Output %zu written to %s (size=%lu bytes)\n", GetTime().c_str(), i, output_file.c_str(), | ||||||||
| 185 | + data_size); | ||||||||
| 186 | + } | ||||||||
| 187 | + | ||||||||
| 188 | + printf("%s - INFO: Finalizing\n", GetTime().c_str()); | ||||||||
| 189 | + delete session; | ||||||||
| 190 | + ret = ge::GEFinalize(); | ||||||||
| 191 | + if (ret != SUCCESS) { | ||||||||
| 192 | + printf("%s - ERROR: GE finalize failed\n", GetTime().c_str()); | ||||||||
| 193 | + return FAILED; | ||||||||
| 194 | + } | ||||||||
| 195 | + printf("%s - INFO: Done\n", GetTime().c_str()); | ||||||||
| 196 | + return SUCCESS; | ||||||||
| 197 | +} | ||||||||
| @@ -0,0 +1,30 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + * | ||
| 10 | + * Generated By CANNBot | ||
| 11 | + */ | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | +using namespace ge; | ||
| 17 | + | ||
| 18 | +namespace ops { | ||
| 19 | + | ||
| 20 | +static ge::graphStatus InferDataTypeForBesselI1e(gert::InferDataTypeContext* context) | ||
| 21 | +{ | ||
| 22 | + OP_LOGI(context->GetNodeName(), "Begin InferDataTypeForBesselI1e"); | ||
| 23 | + const ge::DataType xDataType = context->GetInputDataType(0); | ||
| 24 | + context->SetOutputDataType(0, xDataType); | ||
| 25 | + return ge::GRAPH_SUCCESS; | ||
| 26 | +} | ||
| 27 | + | ||
| 28 | +IMPL_OP(BesselI1e).InferDataType(InferDataTypeForBesselI1e); | ||
| 29 | + | ||
| 30 | +} // namespace ops | ||
| @@ -0,0 +1,42 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + * | ||
| 10 | + * Generated By CANNBot | ||
| 11 | + */ | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | +namespace ge { | ||
| 20 | + | ||
| 21 | +/** | ||
| 22 | + * @brief Computes the Bessel i1e function of "x" element-wise. | ||
| 23 | + * Exponentially scaled modified Bessel function of order 1 | ||
| 24 | + * defined as: bessel_i1e(x) = exp(-abs(x)) * bessel_i1(x). | ||
| 25 | + * | ||
| 26 | + * @par Inputs: | ||
| 27 | + * @li x: A tensor of type bfloat16, float16, float32, or float64. | ||
| 28 | + * | ||
| 29 | + * @par Outputs: | ||
| 30 | + * @li y: A tensor. Has the same dtype as "x". | ||
| 31 | + * | ||
| 32 | + * @par Third-party framework compatibility | ||
| 33 | + * Compatible with the TensorFlow operator BesselI1e. | ||
| 34 | + */ | ||
| 35 | +REG_OP(BesselI1e) | ||
| 36 | + .INPUT(x, TensorType({FloatingDataType, DT_BF16})) | ||
| 37 | + .OUTPUT(y, TensorType({FloatingDataType, DT_BF16})) | ||
| 38 | + .OP_END_FACTORY_REG(BesselI1e) | ||
| 39 | + | ||
| 40 | +} // namespace ge | ||
| 41 | + | ||
| 42 | + | ||
The file is empty
| @@ -0,0 +1,13 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + * | ||
| 10 | + * Generated By CANNBot | ||
| 11 | + */ | ||
| 12 | + | ||
| 13 | + | ||
| @@ -0,0 +1,141 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + * | ||
| 10 | + * Generated By CANNBot | ||
| 11 | + */ | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | +namespace optiling { | ||
| 21 | + | ||
| 22 | +using Ops::Base::CeilAlign; | ||
| 23 | +using Ops::Base::CeilDiv; | ||
| 24 | +using Ops::Base::FloorAlign; | ||
| 25 | +using Ops::Base::FloorDiv; | ||
| 26 | +using Ops::Base::GetUbBlockSize; | ||
| 27 | + | ||
| 28 | +constexpr uint32_t WS_SYS_SIZE = 0U; | ||
| 29 | +constexpr size_t WORKSPACE_NUM = 1; | ||
| 30 | +constexpr int64_t INTERNAL_QUEUE_COUNT = 3; | ||
| 31 | +constexpr int64_t DOUBLE_BUFFER_FACTOR = 2; | ||
| 32 | + | ||
| 33 | +static const gert::Shape g_vec_1_shape = {1}; | ||
| 34 | + | ||
| 35 | +static inline const gert::Shape EnsureNotScalar(const gert::Shape& in_shape) | ||
| 36 | +{ | ||
| 37 | + if (in_shape.GetDimNum() == 0) { | ||
| 38 | + return g_vec_1_shape; | ||
| 39 | + } | ||
| 40 | + return in_shape; | ||
| 41 | +} | ||
| 42 | + | ||
| 43 | +static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t* ubSize, int64_t* coreNum) | ||
| 44 | +{ | ||
| 45 | + fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo(); | ||
| 46 | + OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr); | ||
| 47 | + auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr); | ||
| 48 | + *coreNum = ascendcPlatform.GetCoreNumAiv(); | ||
| 49 | + OP_CHECK_IF(*coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED); | ||
| 50 | + ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, *ubSize); | ||
| 51 | + OP_CHECK_IF(*ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED); | ||
| 52 | + return ge::GRAPH_SUCCESS; | ||
| 53 | +} | ||
| 54 | + | ||
| 55 | +static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t* totalIdx, ge::DataType* dataType) | ||
| 56 | +{ | ||
| 57 | + auto inputX = context->GetInputShape(0); | ||
| 58 | + OP_CHECK_NULL_WITH_CONTEXT(context, inputX); | ||
| 59 | + auto inputShapeX = EnsureNotScalar(inputX->GetStorageShape()); | ||
| 60 | + auto outY = context->GetOutputShape(0); | ||
| 61 | + OP_CHECK_NULL_WITH_CONTEXT(context, outY); | ||
| 62 | + auto outShapeY = EnsureNotScalar(outY->GetStorageShape()); | ||
| 63 | + | ||
| 64 | + OP_CHECK_IF(inputShapeX.GetShapeSize() != outShapeY.GetShapeSize(), | ||
| 65 | + OP_LOGE(context, "BesselI1e: input and output shape size mismatch: x=%ld, y=%ld", | ||
| 66 | + inputShapeX.GetShapeSize(), outShapeY.GetShapeSize()), | ||
| 67 | + return ge::GRAPH_FAILED); | ||
| 68 | + | ||
| 69 | + *totalIdx = inputShapeX.GetShapeSize(); | ||
| 70 | + const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT, ge::DT_FLOAT16, ge::DT_BF16}; | ||
| 71 | + auto inputDesc = context->GetInputDesc(0); | ||
| 72 | + OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc); | ||
| 73 | + *dataType = inputDesc->GetDataType(); | ||
| 74 | + OP_CHECK_IF(supportedDtype.count(*dataType) == 0, OP_LOGE(context, "invalid dtype"), return ge::GRAPH_FAILED); | ||
| 75 | + return ge::GRAPH_SUCCESS; | ||
| 76 | +} | ||
| 77 | + | ||
| 78 | +static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context) | ||
| 79 | +{ | ||
| 80 | + size_t* currentWorkspace = context->GetWorkspaceSizes(WORKSPACE_NUM); | ||
| 81 | + OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace); | ||
| 82 | + currentWorkspace[0] = WS_SYS_SIZE; | ||
| 83 | + return ge::GRAPH_SUCCESS; | ||
| 84 | +} | ||
| 85 | + | ||
| 86 | +static ge::graphStatus BesselI1eTilingFunc(gert::TilingContext* context) | ||
| 87 | +{ | ||
| 88 | + OP_LOGI(context->GetNodeName(), "Enter BesselI1eTilingFunc"); | ||
| 89 | + uint64_t ubSize; | ||
| 90 | + int64_t coreNum; | ||
| 91 | + OP_CHECK_IF(GetPlatformInfo(context, &ubSize, &coreNum) != ge::GRAPH_SUCCESS, | ||
| 92 | + OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED); | ||
| 93 | + | ||
| 94 | + int64_t totalIdx; | ||
| 95 | + ge::DataType dataType; | ||
| 96 | + OP_CHECK_IF(GetShapeAttrsInfo(context, &totalIdx, &dataType) != ge::GRAPH_SUCCESS, | ||
| 97 | + OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED); | ||
| 98 | + | ||
| 99 | + OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, OP_LOGE(context, "GetWorkspaceSize error"), | ||
| 100 | + return ge::GRAPH_FAILED); | ||
| 101 | + | ||
| 102 | + BesselI1eTilingData* tiling = context->GetTilingData<BesselI1eTilingData>(); | ||
| 103 | + OP_CHECK_NULL_WITH_CONTEXT(context, tiling); | ||
| 104 | + OP_CHECK_IF(memset_s(tiling, sizeof(BesselI1eTilingData), 0, sizeof(BesselI1eTilingData)) != EOK, | ||
| 105 | + OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED); | ||
| 106 | + | ||
| 107 | + if (totalIdx == 0) { | ||
| 108 | + context->SetBlockDim(1); | ||
| 109 | + ASCENDC_TPL_SEL_PARAM(context, static_cast<uint32_t>(dataType)); | ||
| 110 | + return ge::GRAPH_SUCCESS; | ||
| 111 | + } | ||
| 112 | + | ||
| 113 | + int64_t ubBlockSize = Ops::Base::GetUbBlockSize(context); | ||
| 114 | + int64_t dtypeSize = (dataType == ge::DT_FLOAT) ? static_cast<int64_t>(sizeof(float)) : | ||
| 115 | + static_cast<int64_t>(sizeof(uint16_t)); | ||
| 116 | + tiling->totalNum = totalIdx; | ||
| 117 | + tiling->blockFactor = CeilAlign(CeilDiv(totalIdx, coreNum), ubBlockSize / dtypeSize); | ||
| 118 | + int64_t usedCoreNum = Ops::Base::CeilDiv(totalIdx, tiling->blockFactor); | ||
| 119 | + | ||
| 120 | + int64_t bufferNum = INTERNAL_QUEUE_COUNT * DOUBLE_BUFFER_FACTOR; | ||
| 121 | + tiling->ubFactor = Ops::Base::FloorAlign( | ||
| 122 | + Ops::Base::FloorDiv((static_cast<int64_t>(ubSize) / static_cast<int64_t>(sizeof(float))), bufferNum), | ||
| 123 | + ubBlockSize / static_cast<int64_t>(sizeof(float))); | ||
| 124 | + | ||
| 125 | + context->SetBlockDim(usedCoreNum); | ||
| 126 | + | ||
| 127 | + uint32_t dTypeX = static_cast<uint32_t>(dataType); | ||
| 128 | + ASCENDC_TPL_SEL_PARAM(context, dTypeX); | ||
| 129 | + return ge::GRAPH_SUCCESS; | ||
| 130 | +} | ||
| 131 | + | ||
| 132 | +static ge::graphStatus TilingParseForBesselI1e([[maybe_unused]] gert::TilingParseContext* context) | ||
| 133 | +{ | ||
| 134 | + return ge::GRAPH_SUCCESS; | ||
| 135 | +} | ||
| 136 | + | ||
| 137 | +struct BesselI1eCompileInfo {}; | ||
| 138 | + | ||
| 139 | +IMPL_OP_OPTILING(BesselI1e).Tiling(BesselI1eTilingFunc).TilingParse<BesselI1eCompileInfo>(TilingParseForBesselI1e); | ||
| 140 | + | ||
| 141 | +} // namespace optiling | ||
| @@ -0,0 +1,47 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + * | ||
| 10 | + * Generated By CANNBot | ||
| 11 | + */ | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | +namespace ops { | ||
| 16 | +class BesselI1e : public OpDef { | ||
| 17 | +public: | ||
| 18 | + explicit BesselI1e(const char* name) : OpDef(name) | ||
| 19 | + { | ||
| 20 | + this->Input("x") | ||
| 21 | + .ParamType(REQUIRED) | ||
| 22 | + .DataType({ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT}) | ||
| 23 | + .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 24 | + .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 25 | + .AutoContiguous(); | ||
| 26 | + this->Output("y") | ||
| 27 | + .ParamType(REQUIRED) | ||
| 28 | + .DataType({ge::DT_BF16, ge::DT_FLOAT16, ge::DT_FLOAT}) | ||
| 29 | + .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 30 | + .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 31 | + .AutoContiguous(); | ||
| 32 | + | ||
| 33 | + OpAICoreConfig aiCoreConfig; | ||
| 34 | + aiCoreConfig.DynamicCompileStaticFlag(true) | ||
| 35 | + .DynamicFormatFlag(false) | ||
| 36 | + .DynamicRankSupportFlag(true) | ||
| 37 | + .DynamicShapeSupportFlag(true) | ||
| 38 | + .NeedCheckSupportFlag(false) | ||
| 39 | + .PrecisionReduceFlag(true) | ||
| 40 | + .ExtendCfgInfo("opFile.value", "bessel_i1e_apt"); | ||
| 41 | + this->AICore().AddConfig("ascend910b", aiCoreConfig); | ||
| 42 | + this->AICore().AddConfig("ascend910_93", aiCoreConfig); | ||
| 43 | + this->AICore().AddConfig("ascend950", aiCoreConfig); | ||
| 44 | + } | ||
| 45 | +}; | ||
| 46 | +OP_ADD(BesselI1e); | ||
| 47 | +} // namespace ops | ||
| @@ -0,0 +1,36 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + * | ||
| 10 | + * Generated By CANNBot | ||
| 11 | + */ | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | +using namespace ge; | ||
| 18 | + | ||
| 19 | +namespace ops { | ||
| 20 | + | ||
| 21 | +static ge::graphStatus InferShape4BesselI1e(gert::InferShapeContext* context) | ||
| 22 | +{ | ||
| 23 | + const gert::Shape* input_shape = context->GetInputShape(0); | ||
| 24 | + OP_CHECK_NULL_WITH_CONTEXT(context, input_shape); | ||
| 25 | + | ||
| 26 | + gert::Shape* output_shape = context->GetOutputShape(0); | ||
| 27 | + OP_CHECK_NULL_WITH_CONTEXT(context, output_shape); | ||
| 28 | + | ||
| 29 | + *output_shape = *input_shape; | ||
| 30 | + | ||
| 31 | + return ge::GRAPH_SUCCESS; | ||
| 32 | +} | ||
| 33 | + | ||
| 34 | +IMPL_OP_INFERSHAPE(BesselI1e).InferShape(InferShape4BesselI1e); | ||
| 35 | + | ||
| 36 | +} // namespace ops | ||
| @@ -0,0 +1,229 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + * | ||
| 10 | + * Generated By CANNBot | ||
| 11 | + */ | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | + | ||
| 20 | + | ||
| 21 | +namespace NsBesselI1e { | ||
| 22 | + | ||
| 23 | +using namespace AscendC; | ||
| 24 | + | ||
| 25 | +constexpr float SEGMENT_POINT = 3.75f; | ||
| 26 | +constexpr float INV_SEGMENT = 0.26666666666666666f; | ||
| 27 | +constexpr float QUARTER = 0.25f; | ||
| 28 | + | ||
| 29 | +constexpr float EXP_COEFF_0 = 1.0f; | ||
| 30 | +constexpr float EXP_COEFF_1 = 1.0f; | ||
| 31 | +constexpr float EXP_COEFF_2 = 0.5f; | ||
| 32 | +constexpr float EXP_COEFF_3 = 0.16666667f; | ||
| 33 | +constexpr float EXP_COEFF_4 = 0.04166667f; | ||
| 34 | +constexpr float EXP_COEFF_5 = 0.00833333f; | ||
| 35 | +constexpr float EXP_COEFF_6 = 0.00138889f; | ||
| 36 | +constexpr float EXP_COEFF_7 = 0.00019841f; | ||
| 37 | +constexpr float EXP_COEFF_8 = 0.00002480f; | ||
| 38 | + | ||
| 39 | +constexpr float itrBefore[7] = {0.5f, 0.87890594f, 0.51498869f, 0.15084934f, 0.02658773f, 0.00301532f, 0.00032411f}; | ||
| 40 | + | ||
| 41 | +constexpr float itrAfter[9] = {0.39894228f, -0.03988024f, -0.00362018f, 0.00163801f, -0.01031555f, | ||
| 42 | + 0.02282967f, -0.02895312f, 0.01787654f, -0.00420059f}; | ||
| 43 | + | ||
| 44 | +template <typename T> | ||
| 45 | +class BesselI1e { | ||
| 46 | + static constexpr int32_t BUFFER_NUM = 2; | ||
| 47 | + | ||
| 48 | +public: | ||
| 49 | + __aicore__ inline BesselI1e(){}; | ||
| 50 | + | ||
| 51 | + __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const BesselI1eTilingData* tilingData); | ||
| 52 | + __aicore__ inline void Process(); | ||
| 53 | + | ||
| 54 | +private: | ||
| 55 | + __aicore__ inline void CopyIn(int64_t progress, int64_t currentNum); | ||
| 56 | + __aicore__ inline void CopyOut(int64_t progress, int64_t currentNum); | ||
| 57 | + __aicore__ inline void Compute(int64_t currentNum); | ||
| 58 | + | ||
| 59 | +private: | ||
| 60 | + TPipe pipe; | ||
| 61 | + TQue<QuePosition::VECIN, BUFFER_NUM> inputQueueX; | ||
| 62 | + TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueueY; | ||
| 63 | + TQue<QuePosition::VECOUT, BUFFER_NUM> tmpQueue1; | ||
| 64 | + | ||
| 65 | + GlobalTensor<T> inputGMX; | ||
| 66 | + GlobalTensor<T> outputGMY; | ||
| 67 | + | ||
| 68 | + int64_t blockLength_ = 0; | ||
| 69 | + int64_t ubLength_ = 0; | ||
| 70 | +}; | ||
| 71 | + | ||
| 72 | +template <typename T> | ||
| 73 | +__aicore__ inline void BesselI1e<T>::Init(GM_ADDR x, GM_ADDR y, const BesselI1eTilingData* tilingData) | ||
| 74 | +{ | ||
| 75 | + int64_t remainderLength = tilingData->totalNum - tilingData->blockFactor * AscendC::GetBlockIdx(); | ||
| 76 | + blockLength_ = (remainderLength > tilingData->blockFactor) ? tilingData->blockFactor : remainderLength; | ||
| 77 | + ubLength_ = tilingData->ubFactor; | ||
| 78 | + | ||
| 79 | + inputGMX.SetGlobalBuffer((__gm__ T*)x + tilingData->blockFactor * AscendC::GetBlockIdx(), blockLength_); | ||
| 80 | + outputGMY.SetGlobalBuffer((__gm__ T*)y + tilingData->blockFactor * AscendC::GetBlockIdx(), blockLength_); | ||
| 81 | + | ||
| 82 | + pipe.InitBuffer(inputQueueX, BUFFER_NUM, ubLength_ * sizeof(float)); | ||
| 83 | + pipe.InitBuffer(outputQueueY, BUFFER_NUM, ubLength_ * sizeof(float)); | ||
| 84 | + pipe.InitBuffer(tmpQueue1, BUFFER_NUM, ubLength_ * sizeof(float)); | ||
| 85 | +} | ||
| 86 | + | ||
| 87 | +template <typename T> | ||
| 88 | +__aicore__ inline void BesselI1e<T>::CopyIn(int64_t progress, int64_t currentNum) | ||
| 89 | +{ | ||
| 90 | + AscendC::LocalTensor<float> xLocal = inputQueueX.template AllocTensor<float>(); | ||
| 91 | + if constexpr (std::is_same_v<T, half>) { | ||
| 92 | + AscendC::LocalTensor<half> tmpHalf = tmpQueue1.template AllocTensor<half>(); | ||
| 93 | + AscendC::DataCopyExtParams copyParams; | ||
| 94 | + copyParams.blockCount = 1; | ||
| 95 | + copyParams.blockLen = currentNum * sizeof(half); | ||
S CopyIn/CopyOut 的尾块 DMA 可能小于 32 字节。当 blockLength_ 不能被 ubFactor 整除时，最后一次迭代 currentNum 可能为 1~31 元素，此时 blockLen = currentNum * sizeof(T) 可能 < 32 字节（fp16 最严重：1 元素 = 2 字节），违反 AscendC DMA 32 字节对齐要求。建议对尾块使用 isPad=true 的 DataCopyPad 配合 DataCopyPadParams 补齐至 32 字节。CopyOut 中同样存在此问题（第 139、150、160 行）。 | |||
| 96 | + copyParams.srcStride = 0; | ||
| 97 | + copyParams.dstStride = 0; | ||
| 98 | + copyParams.rsv = 0; | ||
| 99 | + AscendC::DataCopyPad(tmpHalf, inputGMX[progress * ubLength_], copyParams, {false, 0, 0, 0}); | ||
| 100 | + for (int64_t i = 0; i < currentNum; i++) { | ||
| 101 | + xLocal.SetValue(i, static_cast<float>(tmpHalf.GetValue(i))); | ||
| 102 | + } | ||
| 103 | + tmpQueue1.FreeTensor(tmpHalf); | ||
| 104 | + } else if constexpr (std::is_same_v<T, bfloat16_t>) { | ||
| 105 | + AscendC::LocalTensor<bfloat16_t> tmpBf16 = tmpQueue1.template AllocTensor<bfloat16_t>(); | ||
| 106 | + AscendC::DataCopyExtParams copyParams; | ||
| 107 | + copyParams.blockCount = 1; | ||
| 108 | + copyParams.blockLen = currentNum * sizeof(bfloat16_t); | ||
| 109 | + copyParams.srcStride = 0; | ||
| 110 | + copyParams.dstStride = 0; | ||
| 111 | + copyParams.rsv = 0; | ||
| 112 | + AscendC::DataCopyPad(tmpBf16, inputGMX[progress * ubLength_], copyParams, {false, 0, 0, 0}); | ||
| 113 | + AscendC::Cast(xLocal, tmpBf16, AscendC::RoundMode::CAST_NONE, currentNum); | ||
| 114 | + tmpQueue1.FreeTensor(tmpBf16); | ||
| 115 | + } else { | ||
| 116 | + AscendC::DataCopyExtParams copyParams; | ||
| 117 | + copyParams.blockCount = 1; | ||
| 118 | + copyParams.blockLen = currentNum * sizeof(float); | ||
| 119 | + copyParams.srcStride = 0; | ||
| 120 | + copyParams.dstStride = 0; | ||
| 121 | + copyParams.rsv = 0; | ||
| 122 | + AscendC::DataCopyPad(xLocal, inputGMX[progress * ubLength_], copyParams, {false, 0, 0, 0}); | ||
| 123 | + } | ||
| 124 | + inputQueueX.EnQue(xLocal); | ||
| 125 | +} | ||
| 126 | + | ||
| 127 | +template <typename T> | ||
| 128 | +__aicore__ inline void BesselI1e<T>::CopyOut(int64_t progress, int64_t currentNum) | ||
| 129 | +{ | ||
| 130 | + AscendC::LocalTensor<float> yLocal = outputQueueY.template DeQue<float>(); | ||
| 131 | + if constexpr (std::is_same_v<T, half>) { | ||
| 132 | + AscendC::LocalTensor<half> tmpHalf = tmpQueue1.template AllocTensor<half>(); | ||
| 133 | + for (int64_t i = 0; i < currentNum; i++) { | ||
| 134 | + tmpHalf.SetValue(i, static_cast<half>(yLocal.GetValue(i))); | ||
| 135 | + } | ||
| 136 | + AscendC::DataCopyExtParams copyParams; | ||
| 137 | + copyParams.blockCount = 1; | ||
| 138 | + copyParams.blockLen = currentNum * sizeof(half); | ||
| 139 | + copyParams.srcStride = 0; | ||
| 140 | + copyParams.dstStride = 0; | ||
| 141 | + copyParams.rsv = 0; | ||
| 142 | + AscendC::DataCopyPad(outputGMY[progress * ubLength_], tmpHalf, copyParams); | ||
| 143 | + tmpQueue1.FreeTensor(tmpHalf); | ||
| 144 | + } else if constexpr (std::is_same_v<T, bfloat16_t>) { | ||
| 145 | + AscendC::LocalTensor<bfloat16_t> tmpBf16 = tmpQueue1.template AllocTensor<bfloat16_t>(); | ||
| 146 | + AscendC::Cast(tmpBf16, yLocal, AscendC::RoundMode::CAST_RINT, currentNum); | ||
| 147 | + AscendC::DataCopyExtParams copyParams; | ||
| 148 | + copyParams.blockCount = 1; | ||
| 149 | + copyParams.blockLen = currentNum * sizeof(bfloat16_t); | ||
| 150 | + copyParams.srcStride = 0; | ||
| 151 | + copyParams.dstStride = 0; | ||
| 152 | + copyParams.rsv = 0; | ||
| 153 | + AscendC::DataCopyPad(outputGMY[progress * ubLength_], tmpBf16, copyParams); | ||
| 154 | + tmpQueue1.FreeTensor(tmpBf16); | ||
| 155 | + } else { | ||
| 156 | + AscendC::DataCopyExtParams copyParams; | ||
| 157 | + copyParams.blockCount = 1; | ||
| 158 | + copyParams.blockLen = currentNum * sizeof(float); | ||
| 159 | + copyParams.srcStride = 0; | ||
| 160 | + copyParams.dstStride = 0; | ||
| 161 | + copyParams.rsv = 0; | ||
| 162 | + AscendC::DataCopyPad(outputGMY[progress * ubLength_], yLocal, copyParams); | ||
| 163 | + } | ||
| 164 | + outputQueueY.FreeTensor(yLocal); | ||
| 165 | +} | ||
| 166 | + | ||
| 167 | +template <typename T> | ||
| 168 | +__aicore__ inline void BesselI1e<T>::Compute(int64_t currentNum) | ||
| 169 | +{ | ||
| 170 | + AscendC::LocalTensor<float> xLocal = inputQueueX.template DeQue<float>(); | ||
| 171 | + AscendC::LocalTensor<float> yLocal = outputQueueY.template AllocTensor<float>(); | ||
| 172 | + | ||
| 173 | + for (int64_t i = 0; i < currentNum; i++) { | ||
| 174 | + float x = xLocal.GetValue(i); | ||
| 175 | + float absX = (x >= 0.0f) ? x : -x; | ||
| 176 | + float sign = (x >= 0.0f) ? 1.0f : -1.0f; | ||
| 177 | + float result; | ||
| 178 | + | ||
| 179 | + if (absX < SEGMENT_POINT) { | ||
| 180 | + float t = absX * INV_SEGMENT; | ||
| 181 | + float t2 = t * t; | ||
| 182 | + float poly = itrBefore[6]; | ||
| 183 | + for (int k = 5; k >= 0; k--) { | ||
| 184 | + poly = poly * t2 + itrBefore[k]; | ||
| 185 | + } | ||
| 186 | + float q = absX * QUARTER; | ||
| 187 | + float e = EXP_COEFF_0 - | ||
| 188 | + q * (EXP_COEFF_1 - | ||
| 189 | + q * (EXP_COEFF_2 - | ||
| 190 | + q * (EXP_COEFF_3 - | ||
| 191 | + q * (EXP_COEFF_4 - | ||
| 192 | + q * (EXP_COEFF_5 - | ||
| 193 | + q * (EXP_COEFF_6 - q * (EXP_COEFF_7 - q * EXP_COEFF_8))))))); | ||
| 194 | + e = e * e; | ||
| 195 | + e = e * e; | ||
| 196 | + result = e * absX * poly; | ||
| 197 | + } else { | ||
| 198 | + float t = SEGMENT_POINT / absX; | ||
| 199 | + float poly = itrAfter[8]; | ||
| 200 | + for (int k = 7; k >= 0; k--) { | ||
| 201 | + poly = poly * t + itrAfter[k]; | ||
| 202 | + } | ||
| 203 | + float sqrtX = sqrt(absX); | ||
| 204 | + result = poly / sqrtX; | ||
| 205 | + } | ||
| 206 | + yLocal.SetValue(i, sign * result); | ||
| 207 | + } | ||
| 208 | + | ||
| 209 | + outputQueueY.template EnQue<float>(yLocal); | ||
| 210 | + inputQueueX.FreeTensor(xLocal); | ||
| 211 | +} | ||
| 212 | + | ||
| 213 | +template <typename T> | ||
| 214 | +__aicore__ inline void BesselI1e<T>::Process() | ||
| 215 | +{ | ||
| 216 | + if (blockLength_ <= 0) { | ||
| 217 | + return; | ||
| 218 | + } | ||
| 219 | + int64_t loopCount = (blockLength_ + ubLength_ - 1) / ubLength_; | ||
🟠 High Priority 在 tiling 代码 触发条件：输入为空 tensor（元素个数为 0），如 CSV 测试用例 L0_017/L0_018（shape 修复方向：在 tiling 的 totalIdx==0 分支中将 不准确？ | |||
| 220 | + for (int64_t i = 0; i < loopCount; i++) { | ||
| 221 | + int64_t currentNum = (i == (loopCount - 1)) ? (blockLength_ - ubLength_ * i) : ubLength_; | ||
| 222 | + CopyIn(i, currentNum); | ||
| 223 | + Compute(currentNum); | ||
| 224 | + CopyOut(i, currentNum); | ||
| 225 | + } | ||
| 226 | +} | ||
| 227 | + | ||
| 228 | +} // namespace NsBesselI1e | ||
| 229 | + | ||
| @@ -0,0 +1,21 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + * | ||
| 10 | + * Generated By CANNBot | ||
| 11 | + */ | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | +struct BesselI1eTilingData { | ||
| 17 | + int64_t totalNum = 0; | ||
| 18 | + int64_t blockFactor = 0; | ||
| 19 | + int64_t ubFactor = 0; | ||
| 20 | +}; | ||
| 21 | + | ||
| @@ -0,0 +1,25 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + * | ||
| 10 | + * Generated By CANNBot | ||
| 11 | + */ | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | +ASCENDC_TPL_ARGS_DECL(BesselI1e, | ||
| 19 | + ASCENDC_TPL_DATATYPE_DECL(D_T_X, C_DT_FLOAT, C_DT_FLOAT16, C_DT_BF16, ASCENDC_TPL_INPUT(0))); | ||
| 20 | + | ||
| 21 | +ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT)), | ||
| 22 | + ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_FLOAT16)), | ||
| 23 | + ASCENDC_TPL_ARGS_SEL(ASCENDC_TPL_DATATYPE_SEL(D_T_X, C_DT_BF16)), ); | ||
| 24 | + | ||
| 25 | + | ||
| @@ -0,0 +1,24 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + * | ||
| 10 | + * Generated By CANNBot | ||
| 11 | + */ | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | +template <typename D_T_X> | ||
| 16 | +__global__ __aicore__ void bessel_i1e(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) | ||
| 17 | +{ | ||
| 18 | + KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY); | ||
| 19 | + REGISTER_TILING_DEFAULT(BesselI1eTilingData); | ||
| 20 | + GET_TILING_DATA_WITH_STRUCT(BesselI1eTilingData, tilingData, tiling); | ||
| 21 | + NsBesselI1e::BesselI1e<D_T_X> op; | ||
| 22 | + op.Init(x, y, &tilingData); | ||
| 23 | + op.Process(); | ||
| 24 | +} | ||
| @@ -0,0 +1,8 @@ | |||
| 1 | +cmake_minimum_required(VERSION 3.16.0) | ||
| 2 | +project(bessel_i1e_ut CXX) | ||
| 3 | + | ||
| 4 | +set(CMAKE_CXX_STANDARD 17) | ||
| 5 | +set(CMAKE_CXX_STANDARD_REQUIRED ON) | ||
| 6 | + | ||
| 7 | +add_executable(test_bessel_i1e_ut test_bessel_i1e_ut.cpp) | ||
| 8 | +target_compile_options(test_bessel_i1e_ut PRIVATE -Wall -Wextra -O2) | ||
| @@ -0,0 +1,45 @@ | |||
| 1 | +#!/bin/bash | ||
| 2 | +# Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | +# This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | +# CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | +# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | +# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | +# See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | +# ============================================================================ | ||
| 10 | +set -e | ||
| 11 | + | ||
| 12 | +SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" | ||
| 13 | +BUILD_DIR="${SCRIPT_DIR}/build" | ||
| 14 | + | ||
| 15 | +echo "========================================" | ||
| 16 | +echo "BesselI1e UT Tests" | ||
| 17 | +echo "========================================" | ||
| 18 | + | ||
| 19 | +rm -rf "${BUILD_DIR}" | ||
| 20 | +mkdir -p "${BUILD_DIR}" | ||
| 21 | +cd "${BUILD_DIR}" | ||
| 22 | + | ||
| 23 | +echo "Configuring..." | ||
| 24 | +cmake .. > /dev/null | ||
| 25 | + | ||
| 26 | +echo "Building..." | ||
| 27 | +make -j$(nproc) > /dev/null | ||
| 28 | + | ||
| 29 | +echo "Running..." | ||
| 30 | +echo "" | ||
| 31 | +./test_bessel_i1e_ut | ||
| 32 | +RESULT=$? | ||
| 33 | + | ||
| 34 | +echo "" | ||
| 35 | +if [ $RESULT -eq 0 ]; then | ||
| 36 | + echo "========================================" | ||
| 37 | + echo "UT Result: PASS" | ||
| 38 | + echo "========================================" | ||
| 39 | +else | ||
| 40 | + echo "========================================" | ||
| 41 | + echo "UT Result: FAIL" | ||
| 42 | + echo "========================================" | ||
| 43 | +fi | ||
| 44 | + | ||
| 45 | +exit $RESULT | ||
| @@ -0,0 +1,207 @@ | |||
| 1 | +/** | ||
| 2 | + * Copyright (c) 2026 Huawei Technologies Co., Ltd. | ||
| 3 | + * This program is free software, you can redistribute it and/or modify it under the terms and conditions of | ||
| 4 | + * CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 5 | + * Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 6 | + * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, | ||
| 7 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 8 | + * See LICENSE in the root of the software repository for the full text of the License. | ||
| 9 | + * | ||
| 10 | + * Generated By CANNBot | ||
| 11 | + */ | ||
| 12 | + | ||
| 13 | + | ||
| 14 | + | ||
| 15 | + | ||
| 16 | + | ||
| 17 | + | ||
| 18 | + | ||
| 19 | +static const float kItrBefore[7] = {0.5f, 0.87890594f, 0.51498869f, 0.15084934f, 0.02658773f, 0.00301532f, 0.00032411f}; | ||
| 20 | + | ||
| 21 | +static const float kItrAfter[9] = {0.39894228f, -0.03988024f, -0.00362018f, 0.00163801f, -0.01031555f, | ||
| 22 | + 0.02282967f, -0.02895312f, 0.01787654f, -0.00420059f}; | ||
| 23 | + | ||
| 24 | +static const float kConstLimit = 3.75f; | ||
| 25 | +static const float kInvSegment = 0.26666666666666666f; | ||
| 26 | + | ||
| 27 | +float ComputeI1eScalar(float x) | ||
| 28 | +{ | ||
| 29 | + float absX = fabsf(x); | ||
| 30 | + float sign = (x >= 0.0f) ? 1.0f : -1.0f; | ||
| 31 | + float result; | ||
| 32 | + | ||
| 33 | + if (absX < kConstLimit) { | ||
| 34 | + float t = absX * kInvSegment; | ||
| 35 | + float t2 = t * t; | ||
| 36 | + float poly = kItrBefore[6]; | ||
| 37 | + for (int k = 5; k >= 0; k--) { | ||
| 38 | + poly = poly * t2 + kItrBefore[k]; | ||
| 39 | + } | ||
| 40 | + float q = absX * 0.25f; | ||
| 41 | + float e = 1.0f - | ||
| 42 | + q * (1.0f - q * (0.5f - q * (0.16666667f - | ||
| 43 | + q * (0.04166667f - | ||
| 44 | + q * (0.00833333f - | ||
| 45 | + q * (0.00138889f - q * (0.00019841f - q * 0.00002480f))))))); | ||
| 46 | + e = e * e; | ||
| 47 | + e = e * e; | ||
| 48 | + result = e * absX * poly; | ||
| 49 | + } else { | ||
| 50 | + float t = kConstLimit / absX; | ||
| 51 | + float poly = kItrAfter[8]; | ||
| 52 | + for (int k = 7; k >= 0; k--) { | ||
| 53 | + poly = poly * t + kItrAfter[k]; | ||
| 54 | + } | ||
| 55 | + float sqrtX = sqrtf(absX); | ||
| 56 | + result = poly / sqrtX; | ||
| 57 | + } | ||
| 58 | + return sign * result; | ||
| 59 | +} | ||
| 60 | + | ||
| 61 | +struct TilingData { | ||
| 62 | + int64_t totalNum; | ||
| 63 | + int64_t blockFactor; | ||
| 64 | + int64_t ubFactor; | ||
| 65 | +}; | ||
| 66 | + | ||
| 67 | +void ComputeTiling(int64_t dim0, int32_t typeSize, int64_t ubSize, int64_t coreNum, TilingData& tiling) | ||
| 68 | +{ | ||
| 69 | + constexpr int64_t MIN_TILING_BITS = 32768; | ||
| 70 | + constexpr int64_t ELEM_ALIGN = 512; | ||
| 71 | + constexpr int64_t SINGLE_BUF_COUNT = 5; | ||
| 72 | + constexpr int64_t DOUBLE_BUF_COUNT = 10; | ||
| 73 | + | ||
| 74 | + tiling.totalNum = dim0; | ||
| 75 | + | ||
| 76 | + if (dim0 == 0) { | ||
| 77 | + tiling.blockFactor = 0; | ||
| 78 | + tiling.ubFactor = 0; | ||
| 79 | + return; | ||
| 80 | + } | ||
| 81 | + | ||
| 82 | + int64_t minDtypeBits = typeSize * 8; | ||
| 83 | + int64_t computedCoreNum = (dim0 * minDtypeBits + MIN_TILING_BITS - 1) / MIN_TILING_BITS; | ||
| 84 | + if (computedCoreNum > coreNum) | ||
| 85 | + computedCoreNum = coreNum; | ||
| 86 | + | ||
| 87 | + tiling.blockFactor = ((dim0 + computedCoreNum - 1) / computedCoreNum + ELEM_ALIGN - 1) / ELEM_ALIGN * ELEM_ALIGN; | ||
🟡 Medium Priority 文件 建议：在 不准确？ | |||
| 88 | + | ||
| 89 | + int64_t useDoubleBuffer = (typeSize == 2) ? 1 : ((dim0 > 1024) ? 1 : 0); | ||
| 90 | + int64_t bufferNum = useDoubleBuffer ? DOUBLE_BUF_COUNT : SINGLE_BUF_COUNT; | ||
| 91 | + int64_t bufferDivisor = bufferNum * typeSize; | ||
| 92 | + int64_t maxElemNum = (ubSize * 8) / bufferDivisor; | ||
| 93 | + int64_t alignFactor = 256 * 8 / minDtypeBits; | ||
| 94 | + tiling.ubFactor = (maxElemNum / alignFactor) * alignFactor; | ||
| 95 | +} | ||
| 96 | + | ||
| 97 | +int passed = 0; | ||
| 98 | +int failed = 0; | ||
| 99 | + | ||
| 100 | +void Check(bool cond, const char* name) | ||
| 101 | +{ | ||
| 102 | + if (cond) { | ||
| 103 | + std::cout << " [PASS] " << name << std::endl; | ||
| 104 | + passed++; | ||
| 105 | + } else { | ||
| 106 | + std::cout << " [FAIL] " << name << std::endl; | ||
| 107 | + failed++; | ||
| 108 | + } | ||
| 109 | +} | ||
| 110 | + | ||
| 111 | +void TestGoldenFunction() | ||
| 112 | +{ | ||
| 113 | + std::cout << "\n=== Golden Function Tests ===" << std::endl; | ||
| 114 | + | ||
| 115 | + Check(fabsf(ComputeI1eScalar(0.0f)) < 1e-6f, "i1e(0) = 0"); | ||
| 116 | + Check(fabsf(ComputeI1eScalar(1.0f) - 0.2079104f) < 1e-4f, "i1e(1.0) ≈ 0.2079"); | ||
| 117 | + Check(fabsf(ComputeI1eScalar(-1.0f) + 0.2079104f) < 1e-4f, "i1e(-1.0) ≈ -0.2079 (odd)"); | ||
| 118 | + Check(fabsf(ComputeI1eScalar(2.0f) - 0.2152693f) < 1e-4f, "i1e(2.0) ≈ 0.2153"); | ||
| 119 | + Check(fabsf(ComputeI1eScalar(5.0f) - 0.1639723f) < 1e-4f, "i1e(5.0) ≈ 0.1640"); | ||
| 120 | + Check(fabsf(ComputeI1eScalar(10.0f) - 0.1212627f) < 1e-3f, "i1e(10.0) ≈ 0.1213"); | ||
| 121 | + Check(fabsf(ComputeI1eScalar(100.0f) - 0.0397442f) < 1e-3f, "i1e(100.0) ≈ 0.0397"); | ||
| 122 | + Check(fabsf(ComputeI1eScalar(-5.0f) + 0.1639723f) < 1e-4f, "i1e(-5.0) ≈ -0.1640 (odd)"); | ||
| 123 | + | ||
| 124 | + float seg1 = ComputeI1eScalar(3.74f); | ||
| 125 | + float seg2 = ComputeI1eScalar(3.76f); | ||
| 126 | + Check(fabsf(seg1 - seg2) < 0.01f, "segment continuity at |x|=3.75"); | ||
| 127 | + | ||
| 128 | + Check(ComputeI1eScalar(1000.0f) > 0.0f, "i1e(1000) > 0 (no overflow)"); | ||
| 129 | + Check(ComputeI1eScalar(1000.0f) < 0.1f, "i1e(1000) < 0.1 (asymptotic decay)"); | ||
| 130 | +} | ||
| 131 | + | ||
| 132 | +void TestTilingComputation() | ||
| 133 | +{ | ||
| 134 | + std::cout << "\n=== Tiling Computation Tests ===" << std::endl; | ||
| 135 | + | ||
| 136 | + { | ||
| 137 | + TilingData t; | ||
| 138 | + ComputeTiling(1024, 4, 248 * 1024, 28, t); | ||
| 139 | + Check(t.totalNum == 1024, "FP32 small: totalNum=1024"); | ||
| 140 | + Check(t.blockFactor >= 1024, "FP32 small: blockFactor >= totalNum"); | ||
| 141 | + Check(t.ubFactor > 0, "FP32 small: ubFactor > 0"); | ||
| 142 | + Check(t.ubFactor % 64 == 0, "FP32 small: ubFactor aligned to 64 (256B/4B)"); | ||
| 143 | + } | ||
| 144 | + | ||
| 145 | + { | ||
| 146 | + TilingData t; | ||
| 147 | + ComputeTiling(100000, 4, 248 * 1024, 28, t); | ||
| 148 | + Check(t.totalNum == 100000, "FP32 large: totalNum=100000"); | ||
| 149 | + Check(t.blockFactor % 512 == 0, "FP32 large: blockFactor aligned to 512"); | ||
| 150 | + Check(t.ubFactor > 0, "FP32 large: ubFactor > 0"); | ||
| 151 | + } | ||
| 152 | + | ||
| 153 | + { | ||
| 154 | + TilingData t; | ||
| 155 | + ComputeTiling(512, 2, 248 * 1024, 28, t); | ||
| 156 | + Check(t.totalNum == 512, "FP16 small: totalNum=512"); | ||
| 157 | + Check(t.ubFactor > 0, "FP16 small: ubFactor > 0"); | ||
| 158 | + Check(t.ubFactor % 128 == 0, "FP16 small: ubFactor aligned to 128 (256B/2B)"); | ||
| 159 | + } | ||
| 160 | + | ||
| 161 | + { | ||
| 162 | + TilingData t; | ||
| 163 | + ComputeTiling(1, 4, 248 * 1024, 28, t); | ||
| 164 | + Check(t.totalNum == 1, "FP32 scalar: totalNum=1"); | ||
| 165 | + Check(t.blockFactor == 512, "FP32 scalar: blockFactor=512 (min aligned)"); | ||
| 166 | + } | ||
| 167 | + | ||
| 168 | + { | ||
| 169 | + TilingData t; | ||
| 170 | + ComputeTiling(0, 4, 248 * 1024, 28, t); | ||
| 171 | + Check(t.totalNum == 0, "Empty tensor: totalNum=0"); | ||
| 172 | + } | ||
| 173 | +} | ||
| 174 | + | ||
| 175 | +void TestEdgeCases() | ||
| 176 | +{ | ||
| 177 | + std::cout << "\n=== Edge Case Tests ===" << std::endl; | ||
| 178 | + | ||
| 179 | + float nan_result = ComputeI1eScalar(std::nanf("")); | ||
| 180 | + Check(std::isnan(nan_result), "i1e(NaN) = NaN"); | ||
| 181 | + | ||
| 182 | + float inf_result = ComputeI1eScalar(INFINITY); | ||
| 183 | + Check(inf_result >= 0.0f && inf_result < 0.01f, "i1e(+Inf) ≈ 0"); | ||
| 184 | + | ||
| 185 | + float neg_inf_result = ComputeI1eScalar(-INFINITY); | ||
| 186 | + Check(neg_inf_result <= 0.0f && neg_inf_result > -0.01f, "i1e(-Inf) ≈ 0"); | ||
| 187 | + | ||
| 188 | + float small = ComputeI1eScalar(0.001f); | ||
| 189 | + Check(fabsf(small - 0.0005f) < 0.001f, "i1e(0.001) ≈ 0.0005 (linear region)"); | ||
| 190 | +} | ||
| 191 | + | ||
| 192 | +int main() | ||
| 193 | +{ | ||
| 194 | + std::cout << "========================================" << std::endl; | ||
| 195 | + std::cout << "BesselI1e Unit Tests" << std::endl; | ||
| 196 | + std::cout << "========================================" << std::endl; | ||
| 197 | + | ||
| 198 | + TestGoldenFunction(); | ||
| 199 | + TestTilingComputation(); | ||
| 200 | + TestEdgeCases(); | ||
| 201 | + | ||
| 202 | + std::cout << "\n========================================" << std::endl; | ||
| 203 | + std::cout << "Results: " << passed << " passed, " << failed << " failed" << std::endl; | ||
| 204 | + std::cout << "========================================" << std::endl; | ||
| 205 | + | ||
| 206 | + return failed == 0 ? 0 : 1; | ||
| 207 | +} | ||
| @@ -71,6 +71,18 @@ do | |||
| 71 | done | 71 | done |
| 72 | fi | 72 | fi |
| 73 | 73 | ||
| 74 | + if [[ ! "$file_path" == "experimental/"* ]]; then | ||
| 75 | + for dir in "${builtin_dirs[@]}" | ||
| 76 | + do | ||
| 77 | + if [[ "$file_path" == *"/$dir/"*"/arch35/"* ]]; then | ||
| 78 | + if [[ ! " ${builtin_ops_name[@]} " =~ " $dir " ]]; then | ||
| 79 | + builtin_ops_name+=("$dir") | ||
| 80 | + fi | ||
| 81 | + break | ||
| 82 | + fi | ||
| 83 | + done | ||
| 84 | + fi | ||
| 85 | + | ||
| 74 | for dir in "${experimental_dirs[@]}" | 86 | for dir in "${experimental_dirs[@]}" |
| 75 | do | 87 | do |
| 76 | if [[ "$file_path" == "experimental/"*"/$dir/"*"/arch35/"* ]]; then | 88 | if [[ "$file_path" == "experimental/"*"/$dir/"*"/arch35/"* ]]; then |
| @@ -119,14 +131,14 @@ execute_run_file() { | |||
| 119 | 131 | ||
| 120 | if [ ${#builtin_ops_name[@]} -gt 0 ]; then | 132 | if [ ${#builtin_ops_name[@]} -gt 0 ]; then |
| 121 | builtin_ops_str=$(IFS=,; echo "${builtin_ops_name[*]}") | 133 | builtin_ops_str=$(IFS=,; echo "${builtin_ops_name[*]}") |
| 122 | - build_cmd="bash build.sh --pkg --ops=$builtin_ops_str --soc=ascend950 -j16 --cann_3rd_lib_path=${ASCEND_3RD_LIB_PATH}" | 134 | + build_cmd="bash build.sh --pkg --ops=${builtin_ops_str} --soc=ascend950 -j16 --cann_3rd_lib_path=${ASCEND_3RD_LIB_PATH}" |
| 123 | run_build_command "$build_cmd" | 135 | run_build_command "$build_cmd" |
| 124 | execute_run_file "custom" | 136 | execute_run_file "custom" |
| 125 | fi | 137 | fi |
| 126 | 138 | ||
| 127 | if [ ${#experimental_ops_name[@]} -gt 0 ]; then | 139 | if [ ${#experimental_ops_name[@]} -gt 0 ]; then |
| 128 | experimental_ops_str=$(IFS=,; echo "${experimental_ops_name[*]}") | 140 | experimental_ops_str=$(IFS=,; echo "${experimental_ops_name[*]}") |
| 129 | - build_cmd="bash build.sh --pkg --experimental --ops=$experimental_ops_str --soc=ascend950 -j16 --cann_3rd_lib_path=${ASCEND_3RD_LIB_PATH}" | 141 | + build_cmd="bash build.sh --pkg --experimental --ops=${experimental_ops_str} --soc=ascend950 -j16 --cann_3rd_lib_path=${ASCEND_3RD_LIB_PATH}" |
| 130 | run_build_command "$build_cmd" | 142 | run_build_command "$build_cmd" |
| 131 | execute_run_file "custom" | 143 | execute_run_file "custom" |
| 132 | fi | 144 | fi |
缺少CANNBot标签，排查修改。