已合并
Square算子AscendC实现贡献 #291
liwen创建于 2025年11月24日
Square算子AscendC实现贡献 #291
已合并
共 18 个文件变更+1308-0
| @@ -0,0 +1,12 @@ | |||
| 1 | +# ---------------------------------------------------------------------------- | ||
| 2 | +# This program is free software, you can redistribute it and/or modify it. | ||
| 3 | +# Copyright (c) 2025 Huawei Technologies Co., Ltd. | ||
| 4 | +# This file is a part of the CANN Open Software. | ||
| 5 | +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 6 | +# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 7 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | ||
| 8 | +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 9 | +# See LICENSE in the root of the software repository for the full text of the License. | ||
| 10 | +# ---------------------------------------------------------------------------- | ||
| 11 | + | ||
| 12 | +add_all_modules_sources(OPTYPE square_v2 ACLNNTYPE aclnn) | ||
| @@ -0,0 +1,170 @@ | |||
| 1 | +/** | ||
| 2 | + * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT) | ||
| 3 | + * and is contributed to the CANN Open Software. | ||
| 4 | + * | ||
| 5 | + * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT). | ||
| 6 | + * All Rights Reserved. | ||
| 7 | + * | ||
| 8 | + * Authors (accounts): | ||
| 9 | + * - Li Wen <@liwenkkklll> | ||
| 10 | + * - Su Tonghua <@sutonghua> | ||
| 11 | + * | ||
| 12 | + * This program is free software: you can redistribute it and/or modify it. | ||
| 13 | + * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 14 | + * You may not use this file except in compliance with the License. | ||
| 15 | + * See the LICENSE file at the root of the repository for the full text of the License. | ||
| 16 | + * | ||
| 17 | + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED, | ||
| 18 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 19 | + */ | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + do { \ | ||
| 28 | + if (!(cond)) { \ | ||
| 29 | + return_expr; \ | ||
| 30 | + } \ | ||
| 31 | + } while (0) | ||
| 32 | + | ||
| 33 | + | ||
| 34 | + do { \ | ||
| 35 | + printf(message, ##__VA_ARGS__); \ | ||
| 36 | + } while (0) | ||
| 37 | + | ||
| 38 | + int64_t GetShapeSize(const std::vector<int64_t>& shape) | ||
| 39 | + { | ||
| 40 | + int64_t shapeSize = 1; | ||
| 41 | + for (auto i : shape) { | ||
| 42 | + shapeSize *= i; | ||
| 43 | + } | ||
| 44 | + return shapeSize; | ||
| 45 | + } | ||
| 46 | + | ||
| 47 | + void PrintOutResult(std::vector<int64_t>& shape, void** deviceAddr) | ||
| 48 | + { | ||
| 49 | + auto size = GetShapeSize(shape); | ||
| 50 | + std::vector<float> resultData(size, 0); | ||
| 51 | + auto ret = aclrtMemcpy( | ||
| 52 | + resultData.data(), resultData.size() * sizeof(resultData[0]), *deviceAddr, size * sizeof(resultData[0]), | ||
| 53 | + ACL_MEMCPY_DEVICE_TO_HOST); | ||
| 54 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return); | ||
| 55 | + for (int64_t i = 0; i < size; i++) { | ||
| 56 | + LOG_PRINT("mean result[%ld] is: %f\n", i, resultData[i]); | ||
| 57 | + } | ||
| 58 | + } | ||
| 59 | + | ||
| 60 | + int Init(int32_t deviceId, aclrtStream* stream) | ||
| 61 | + { | ||
| 62 | + // 固定写法,初始化 | ||
| 63 | + auto ret = aclInit(nullptr); | ||
| 64 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); | ||
| 65 | + ret = aclrtSetDevice(deviceId); | ||
| 66 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); | ||
| 67 | + ret = aclrtCreateStream(stream); | ||
| 68 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); | ||
| 69 | + return 0; | ||
| 70 | + } | ||
| 71 | + | ||
| 72 | + template <typename T> | ||
| 73 | + int CreateAclTensor( | ||
| 74 | + const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType, | ||
| 75 | + aclTensor** tensor) | ||
| 76 | + { | ||
| 77 | + auto size = GetShapeSize(shape) * sizeof(T); | ||
| 78 | + // 2. 申请device侧内存 | ||
| 79 | + auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 80 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret); | ||
| 81 | + // 3. 调用aclrtMemcpy将host侧数据拷贝到device侧内存上 | ||
| 82 | + ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE); | ||
| 83 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret); | ||
| 84 | + | ||
| 85 | + // 计算连续tensor的strides | ||
| 86 | + std::vector<int64_t> strides(shape.size(), 1); | ||
| 87 | + for (int64_t i = shape.size() - 2; i >= 0; i--) { | ||
| 88 | + strides[i] = shape[i + 1] * strides[i + 1]; | ||
| 89 | + } | ||
| 90 | + | ||
| 91 | + // 调用aclCreateTensor接口创建aclTensor | ||
| 92 | + *tensor = aclCreateTensor( | ||
| 93 | + shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), | ||
| 94 | + *deviceAddr); | ||
| 95 | + return 0; | ||
| 96 | + } | ||
| 97 | + | ||
| 98 | + int main() | ||
| 99 | + { | ||
| 100 | + // 1. 调用acl进行device/stream初始化 | ||
| 101 | + int32_t deviceId = 0; | ||
| 102 | + aclrtStream stream; | ||
| 103 | + auto ret = Init(deviceId, &stream); | ||
| 104 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret); | ||
| 105 | + | ||
| 106 | + // 2. 构造输入与输出,需要根据API的接口自定义构造 | ||
| 107 | + aclTensor* selfX = nullptr; | ||
| 108 | + void* selfXDeviceAddr = nullptr; | ||
| 109 | + std::vector<int64_t> selfXShape = {9}; | ||
| 110 | + std::vector<float> selfXHostData(9, 2.5f); | ||
| 111 | + ret = CreateAclTensor(selfXHostData, selfXShape, &selfXDeviceAddr, aclDataType::ACL_FLOAT, &selfX); | ||
| 112 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 113 | + | ||
| 114 | + aclTensor* out = nullptr; | ||
| 115 | + void* outDeviceAddr = nullptr; | ||
| 116 | + std::vector<int64_t> outShape = {9}; | ||
| 117 | + std::vector<float> outHostData(9, 0); | ||
| 118 | + ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out); | ||
| 119 | + CHECK_RET(ret == ACL_SUCCESS, return ret); | ||
| 120 | + | ||
| 121 | + // 3. 调用CANN算子库API,需要修改为具体的Api名称 | ||
| 122 | + uint64_t workspaceSize = 0; | ||
| 123 | + aclOpExecutor* executor; | ||
| 124 | + | ||
| 125 | + // 4. 调用aclnnSquareV2第一段接口 | ||
| 126 | + ret = aclnnSquareV2GetWorkspaceSize(selfX, out, &workspaceSize, &executor); | ||
| 127 | + // CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSquareV2GetWorkspaceSize failed. ERROR: %d\n", ret); return ret); | ||
| 128 | + if (ret != ACL_SUCCESS) { | ||
| 129 | + const char* errMsg = aclGetRecentErrMsg(); // ← 关键一行 | ||
| 130 | + LOG_PRINT("[ERROR] aclnnSquareV2GetWorkspaceSize failed, ret = %d\n", ret); | ||
| 131 | + LOG_PRINT(" recent AscendCL msg: %s\n", | ||
| 132 | + errMsg ? errMsg : "no additional message"); | ||
| 133 | + return ret; | ||
| 134 | + } | ||
| 135 | + // 根据第一段接口计算出的workspaceSize申请device内存 | ||
| 136 | + void* workspaceAddr = nullptr; | ||
| 137 | + if (workspaceSize > static_cast<uint64_t>(0)) { | ||
| 138 | + ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); | ||
| 139 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret); | ||
| 140 | + } | ||
| 141 | + | ||
| 142 | + // 5. 调用aclnnSquareV2第二段接口 | ||
| 143 | + ret = aclnnSquareV2(workspaceAddr, workspaceSize, executor, stream); | ||
| 144 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSquareV2 failed. ERROR: %d\n", ret); return ret); | ||
| 145 | + | ||
| 146 | + // 6. (固定写法)同步等待任务执行结束 | ||
| 147 | + ret = aclrtSynchronizeStream(stream); | ||
| 148 | + CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); | ||
| 149 | + | ||
| 150 | + // 5. 获取输出的值,将device侧内存上的结果拷贝至host侧,需要根据具体API的接口定义修改 | ||
| 151 | + PrintOutResult(outShape, &outDeviceAddr); | ||
| 152 | + | ||
| 153 | + // 7. 释放aclTensor,需要根据具体API的接口定义修改 | ||
| 154 | + aclDestroyTensor(selfX); | ||
| 155 | + aclDestroyTensor(out); | ||
| 156 | + | ||
| 157 | + // 8. 释放device资源 | ||
| 158 | + aclrtFree(selfXDeviceAddr); | ||
| 159 | + aclrtFree(outDeviceAddr); | ||
| 160 | + if (workspaceSize > static_cast<uint64_t>(0)) { | ||
| 161 | + aclrtFree(workspaceAddr); | ||
| 162 | + } | ||
| 163 | + aclrtDestroyStream(stream); | ||
| 164 | + aclrtResetDevice(deviceId); | ||
| 165 | + | ||
| 166 | + // 9. acl去初始化 | ||
| 167 | + aclFinalize(); | ||
| 168 | + | ||
| 169 | + return 0; | ||
| 170 | + } | ||
| @@ -0,0 +1,325 @@ | |||
| 1 | +/** | ||
| 2 | + * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT) | ||
| 3 | + * and is contributed to the CANN Open Software. | ||
| 4 | + * | ||
| 5 | + * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT). | ||
| 6 | + * All Rights Reserved. | ||
| 7 | + * | ||
| 8 | + * Authors (accounts): | ||
| 9 | + * - Li Wen <@liwenkkklll> | ||
| 10 | + * - Su Tonghua <@sutonghua> | ||
| 11 | + * | ||
| 12 | + * This program is free software: you can redistribute it and/or modify it. | ||
| 13 | + * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 14 | + * You may not use this file except in compliance with the License. | ||
| 15 | + * See the LICENSE file at the root of the repository for the full text of the License. | ||
| 16 | + * | ||
| 17 | + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED, | ||
| 18 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 19 | + */ | ||
| 20 | + | ||
| 21 | + | ||
| 22 | + | ||
| 23 | + | ||
| 24 | + | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | + | ||
| 34 | + | ||
| 35 | + | ||
| 36 | + | ||
| 37 | + | ||
| 38 | + | ||
| 39 | + | ||
| 40 | + | ||
| 41 | + | ||
| 42 | + | ||
| 43 | + | ||
| 44 | + | ||
| 45 | +using namespace ge; | ||
| 46 | +using std::map; | ||
| 47 | +using std::string; | ||
| 48 | +using std::vector; | ||
| 49 | + | ||
| 50 | + vector<int64_t> placeholder##intputIndex##_shape = inputShape; \ | ||
| 51 | + auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \ | ||
| 52 | + TensorDesc placeholder##intputIndex##_desc = \ | ||
| 53 | + TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \ | ||
| 54 | + placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \ | ||
| 55 | + placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \ | ||
| 56 | + Tensor tensor_placeholder##intputIndex; \ | ||
| 57 | + if (intputDtype == DT_FLOAT || intputDtype == DT_FLOAT16) { \ | ||
| 58 | + ret = GenOnesDataFloat32(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \ | ||
| 59 | + placeholder##intputIndex##_desc, static_cast<float>(value)); \ | ||
| 60 | + } else if(intputDtype == DT_INT32 || intputDtype == DT_INT16){ \ | ||
| 61 | + ret = GenOnesData(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \ | ||
| 62 | + placeholder##intputIndex##_desc, intputDtype, static_cast<int>(value)); \ | ||
| 63 | + } \ | ||
| 64 | + if (ret != SUCCESS) { \ | ||
| 65 | + printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \ | ||
| 66 | + return FAILED; \ | ||
| 67 | + } \ | ||
| 68 | + placeholder##intputIndex.update_input_desc_x(placeholder##intputIndex##_desc); \ | ||
| 69 | + input.push_back(tensor_placeholder##intputIndex); \ | ||
| 70 | + graph.AddOp(placeholder##intputIndex); \ | ||
| 71 | + add1.set_input_##intputName(placeholder##intputIndex); \ | ||
| 72 | + inputs.push_back(placeholder##intputIndex); | ||
| 73 | + | ||
| 74 | + | ||
| 75 | + vector<int64_t> placeholder##intputIndex##_shape = inputShape; \ | ||
| 76 | + auto placeholder##intputIndex = op::Const("placeholder" + intputIndex); \ | ||
| 77 | + TensorDesc placeholder##intputIndex##_desc = \ | ||
| 78 | + TensorDesc(ge::Shape(placeholder##intputIndex##_shape), FORMAT_ND, intputDtype); \ | ||
| 79 | + placeholder##intputIndex##_desc.SetPlacement(ge::kPlacementHost); \ | ||
| 80 | + placeholder##intputIndex##_desc.SetFormat(FORMAT_ND); \ | ||
| 81 | + Tensor tensor_placeholder##intputIndex; \ | ||
| 82 | + if (intputDtype == DT_FLOAT || intputDtype == DT_FLOAT16) { \ | ||
| 83 | + ret = GenOnesDataFloat32(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \ | ||
| 84 | + placeholder##intputIndex##_desc, static_cast<float>(value)); \ | ||
| 85 | + } else if(intputDtype == DT_INT32 || intputDtype == DT_INT16){ \ | ||
| 86 | + ret = GenOnesData(placeholder##intputIndex##_shape, tensor_placeholder##intputIndex, \ | ||
| 87 | + placeholder##intputIndex##_desc, intputDtype, static_cast<int>(value)); \ | ||
| 88 | + } \ | ||
| 89 | + if (ret != SUCCESS) { \ | ||
| 90 | + printf("%s - ERROR - [XIR]: Generate input data failed\n", GetTime().c_str()); \ | ||
| 91 | + return FAILED; \ | ||
| 92 | + } \ | ||
| 93 | + placeholder##intputIndex.SetAttr("value", tensor_placeholder##intputIndex); \ | ||
| 94 | + placeholder##intputIndex.update_output_desc_y(placeholder##intputIndex##_desc); \ | ||
| 95 | + graph.AddOp(placeholder##intputIndex); \ | ||
| 96 | + add1.set_input_##intputName(placeholder##intputIndex); \ | ||
| 97 | + add1.update_input_desc_##intputName(placeholder##intputIndex##_desc); \ | ||
| 98 | + inputs.push_back(placeholder##intputIndex); | ||
| 99 | + | ||
| 100 | + | ||
| 101 | + TensorDesc outputName##outputIndex##_desc = TensorDesc(ge::Shape(outputShape), FORMAT_ND, outputDtype); \ | ||
| 102 | + add1.update_output_desc_##outputName(outputName##outputIndex##_desc); | ||
| 103 | + //新加 | ||
| 104 | + | ||
| 105 | + do { \ | ||
| 106 | + printf(message, ##__VA_ARGS__); \ | ||
| 107 | + } while (0) | ||
| 108 | + | ||
| 109 | +string GetTime() | ||
| 110 | +{ | ||
| 111 | + time_t timep; | ||
| 112 | + time(&timep); | ||
| 113 | + char tmp[64]; | ||
| 114 | + strftime(tmp, sizeof(tmp), "%Y-%m-%d %H:%M:%S,000", localtime(&timep)); | ||
| 115 | + return tmp; | ||
| 116 | +} | ||
| 117 | + | ||
| 118 | +uint32_t GetDataTypeSize(DataType dt) | ||
| 119 | +{ | ||
| 120 | + uint32_t dilation = 1; | ||
| 121 | + uint32_t oneByte = 1; | ||
| 122 | + uint32_t twoByte = 2; | ||
| 123 | + uint32_t fourByte = 4; | ||
| 124 | + uint32_t eightByte = 8; | ||
| 125 | + | ||
| 126 | + if (dt == ge::DT_FLOAT) { | ||
| 127 | + dilation = fourByte; | ||
| 128 | + } else if (dt == ge::DT_FLOAT16) { | ||
| 129 | + dilation = twoByte; | ||
| 130 | + } else if (dt == ge::DT_BF16) { | ||
| 131 | + dilation = twoByte; | ||
| 132 | + } else if (dt == ge::DT_INT16) { | ||
| 133 | + dilation = twoByte; | ||
| 134 | + } else if (dt == ge::DT_UINT16) { | ||
| 135 | + dilation = twoByte; | ||
| 136 | + } else if (dt == ge::DT_INT32) { | ||
| 137 | + dilation = fourByte; | ||
| 138 | + } else if (dt == ge::DT_UINT32) { | ||
| 139 | + dilation = fourByte; | ||
| 140 | + } else if (dt == ge::DT_INT64) { | ||
| 141 | + dilation = eightByte; | ||
| 142 | + } else if (dt == ge::DT_UINT64) { | ||
| 143 | + dilation = eightByte; | ||
| 144 | + } else if (dt == ge::DT_INT8) { | ||
| 145 | + dilation = oneByte; | ||
| 146 | + } | ||
| 147 | + return dilation; | ||
| 148 | +} | ||
| 149 | + | ||
| 150 | +int32_t GenOnesDataFloat32(vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, float value) | ||
| 151 | +{ | ||
| 152 | + input_tensor_desc.SetRealDimCnt(shapes.size()); | ||
| 153 | + size_t size = 1; | ||
| 154 | + for (uint32_t i = 0; i < shapes.size(); i++) { | ||
| 155 | + size *= shapes[i]; | ||
| 156 | + } | ||
| 157 | + uint32_t byteSizeFloat32 = 4; | ||
| 158 | + uint32_t data_len = size * byteSizeFloat32; | ||
| 159 | + float* pData = new (std::nothrow) float[size]; | ||
| 160 | + | ||
| 161 | + for (size_t i = 0; i < size; ++i) { | ||
| 162 | + *(pData + i) = value; | ||
| 163 | + } | ||
| 164 | + input_tensor = Tensor(input_tensor_desc, (uint8_t*)pData, data_len); | ||
| 165 | + return SUCCESS; | ||
| 166 | +} | ||
| 167 | + | ||
| 168 | +int32_t GenOnesData( | ||
| 169 | + vector<int64_t> shapes, Tensor& input_tensor, TensorDesc& input_tensor_desc, DataType data_type, int value) | ||
| 170 | +{ | ||
| 171 | + input_tensor_desc.SetRealDimCnt(shapes.size()); | ||
| 172 | + size_t size = 1; | ||
| 173 | + for (uint32_t i = 0; i < shapes.size(); i++) { | ||
| 174 | + size *= shapes[i]; | ||
| 175 | + } | ||
| 176 | + uint32_t data_len = size * GetDataTypeSize(data_type); | ||
| 177 | + int32_t* pData = new (std::nothrow) int32_t[size]; | ||
| 178 | + for (uint32_t i = 0; i < size; ++i) { | ||
| 179 | + *(pData + i) = value; | ||
| 180 | + } | ||
| 181 | + input_tensor = Tensor(input_tensor_desc, reinterpret_cast<uint8_t*>(pData), data_len); | ||
| 182 | + return SUCCESS; | ||
| 183 | +} | ||
| 184 | + | ||
| 185 | +int32_t WriteDataToFile(string bin_file, uint64_t data_size, uint8_t* inputData) | ||
| 186 | +{ | ||
| 187 | + FILE* fp; | ||
| 188 | + fp = fopen(bin_file.c_str(), "w"); | ||
| 189 | + fwrite(inputData, sizeof(uint8_t), data_size, fp); | ||
| 190 | + fclose(fp); | ||
| 191 | + return SUCCESS; | ||
| 192 | +} | ||
| 193 | + | ||
| 194 | +int CreateOppInGraph( | ||
| 195 | + DataType inDtype, std::vector<ge::Tensor>& input, std::vector<Operator>& inputs, std::vector<Operator>& outputs, | ||
| 196 | + Graph& graph) | ||
| 197 | +{ | ||
| 198 | + Status ret = SUCCESS; | ||
| 199 | + // 自定义代码:添加单算子定义到图中 | ||
| 200 | + auto add1 = op::SquareV2("add1"); | ||
| 201 | + std::vector<int64_t> xShape = {1, 4, 4, 4}; | ||
| 202 | + ADD_INPUT(1, x, inDtype, xShape,6.1f); | ||
| 203 | + | ||
| 204 | + ADD_OUTPUT(1, z, inDtype, xShape); | ||
| 205 | + | ||
| 206 | + outputs.push_back(add1); | ||
| 207 | + // 添加完毕 | ||
| 208 | + return SUCCESS; | ||
| 209 | +} | ||
| 210 | + | ||
| 211 | +int main(int argc, char* argv[]) | ||
| 212 | +{ | ||
| 213 | + const char* graph_name = "tc_ge_irrun_test"; | ||
| 214 | + Graph graph(graph_name); | ||
| 215 | + std::vector<ge::Tensor> input; | ||
| 216 | + | ||
| 217 | + printf("%s - INFO - [XIR]: Start to initialize ge using ge global options\n", GetTime().c_str()); | ||
| 218 | + std::map<AscendString, AscendString> global_options = {{"ge.exec.deviceId", "0"}, {"ge.graphRunMode", "1"}}; | ||
| 219 | + Status ret = ge::GEInitialize(global_options); | ||
| 220 | + if (ret != SUCCESS) { | ||
| 221 | + printf("%s - INFO - [XIR]: Initialize ge using ge global options failed\n", GetTime().c_str()); | ||
| 222 | + return FAILED; | ||
| 223 | + } | ||
| 224 | + printf("%s - INFO - [XIR]: Initialize ge using ge global options success\n", GetTime().c_str()); | ||
| 225 | + | ||
| 226 | + std::vector<Operator> inputs{}; | ||
| 227 | + std::vector<Operator> outputs{}; | ||
| 228 | + | ||
| 229 | + std::cout << argv[1] << std::endl; | ||
| 230 | + char* endptr; | ||
| 231 | + //修改类型 | ||
| 232 | + // DataType inDtype = DT_INT32; | ||
| 233 | + DataType inDtype = DT_FLOAT; | ||
| 234 | + | ||
| 235 | + std::cout << inDtype << std::endl; | ||
| 236 | + | ||
| 237 | + ret = CreateOppInGraph(inDtype, input, inputs, outputs, graph); | ||
| 238 | + if (ret != SUCCESS) { | ||
| 239 | + printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str()); | ||
| 240 | + return FAILED; | ||
| 241 | + } | ||
| 242 | + | ||
| 243 | + if (!inputs.empty() && !outputs.empty()) { | ||
| 244 | + graph.SetInputs(inputs).SetOutputs(outputs); | ||
| 245 | + } | ||
| 246 | + | ||
| 247 | + std::map<AscendString, AscendString> build_options = { | ||
| 248 | + | ||
| 249 | + }; | ||
| 250 | + printf("%s - INFO - [XIR]: Start to create ir session using build options\n", GetTime().c_str()); | ||
| 251 | + ge::Session* session = new Session(build_options); | ||
| 252 | + | ||
| 253 | + if (session == nullptr) { | ||
| 254 | + printf("%s - ERROR - [XIR]: Create ir session using build options failed\n", GetTime().c_str()); | ||
| 255 | + return FAILED; | ||
| 256 | + } | ||
| 257 | + printf("%s - INFO - [XIR]: Create ir session using build options success\n", GetTime().c_str()); | ||
| 258 | + printf("%s - INFO - [XIR]: Start to add compute graph to ir session\n", GetTime().c_str()); | ||
| 259 | + | ||
| 260 | + std::map<AscendString, AscendString> graph_options = { | ||
| 261 | + | ||
| 262 | + }; | ||
| 263 | + uint32_t graph_id = 0; | ||
| 264 | + ret = session->AddGraph(graph_id, graph, graph_options); | ||
| 265 | + | ||
| 266 | + printf("%s - INFO - [XIR]: Session add ir compute graph to ir session success\n", GetTime().c_str()); | ||
| 267 | + printf("%s - INFO - [XIR]: dump graph to txt\n", GetTime().c_str()); | ||
| 268 | + std::string file_path = "./dump"; | ||
| 269 | + aclgrphDumpGraph(graph, file_path.c_str(), file_path.length()); | ||
| 270 | + printf("%s - INFO - [XIR]: Start to run ir compute graph\n", GetTime().c_str()); | ||
| 271 | + std::vector<ge::Tensor> output; | ||
| 272 | + ret = session->RunGraph(graph_id, input, output); | ||
| 273 | + if (ret != SUCCESS) { | ||
| 274 | + printf("%s - INFO - [XIR]: Run graph failed\n", GetTime().c_str()); | ||
| 275 | + delete session; | ||
| 276 | + GEFinalize(); | ||
| 277 | + return FAILED; | ||
| 278 | + } | ||
| 279 | + printf("%s - INFO - [XIR]: Session run ir compute graph success\n", GetTime().c_str()); | ||
| 280 | + | ||
| 281 | + int input_num = input.size(); | ||
| 282 | + for (int i = 0; i < input_num; i++) { | ||
| 283 | + std::cout << "input " << i << " dtype : " << input[i].GetTensorDesc().GetDataType() << std::endl; | ||
| 284 | + string input_file = "./tc_ge_irrun_test_0008_npu_input_" + std::to_string(i) + ".bin"; | ||
| 285 | + uint8_t* input_data_i = input[i].GetData(); | ||
| 286 | + int64_t input_shape = input[i].GetTensorDesc().GetShape().GetShapeSize(); | ||
| 287 | + std::cout << "this is " << i << "th input, input shape size =" << input_shape << std::endl; | ||
| 288 | + uint32_t data_size = input_shape * GetDataTypeSize(input[i].GetTensorDesc().GetDataType()); | ||
| 289 | + WriteDataToFile((const char*)input_file.c_str(), data_size, input_data_i); | ||
| 290 | + } | ||
| 291 | + | ||
| 292 | + int output_num = output.size(); | ||
| 293 | + for (int i = 0; i < output_num; i++) { | ||
| 294 | + std::cout << "output " << i << " dtype : " << output[i].GetTensorDesc().GetDataType() << std::endl; | ||
| 295 | + string output_file = "./tc_ge_irrun_test_0008_npu_output_" + std::to_string(i) + ".bin"; | ||
| 296 | + uint8_t* output_data_i = output[i].GetData(); | ||
| 297 | + int64_t output_shape = output[i].GetTensorDesc().GetShape().GetShapeSize(); | ||
| 298 | + std::cout << "this is " << i << "th output, output shape size =" << output_shape << std::endl; | ||
| 299 | + uint32_t data_size = output_shape * GetDataTypeSize(output[i].GetTensorDesc().GetDataType()); | ||
| 300 | + WriteDataToFile((const char*)output_file.c_str(), data_size, output_data_i); | ||
| 301 | + //新加3行打印 | ||
| 302 | + float* result = (float*)output_data_i; | ||
| 303 | + // int32_t* result = (int32_t*)output_data_i; | ||
| 304 | + for (int64_t j = 0; j < 8; j++) { | ||
| 305 | + LOG_PRINT("result[%ld] is: %f\n", j, result[j]); | ||
| 306 | + // LOG_PRINT("result[%ld] is: %d\n", j, result[j]); | ||
| 307 | + } | ||
| 308 | + } | ||
| 309 | + | ||
| 310 | + ge::AscendString error_msg = ge::GEGetErrorMsgV2(); | ||
| 311 | + std::string error_str(error_msg.GetString()); | ||
| 312 | + std::cout << "Error message: " << error_str << std::endl; | ||
| 313 | + ge::AscendString warning_msg = ge::GEGetWarningMsgV2(); | ||
| 314 | + std::string warning_str(warning_msg.GetString()); | ||
| 315 | + std::cout << "Warning message: " << warning_str << std::endl; | ||
| 316 | + printf("%s - INFO - [XIR]: Precision is ok\n", GetTime().c_str()); | ||
| 317 | + printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str()); | ||
| 318 | + ret = ge::GEFinalize(); | ||
| 319 | + if (ret != SUCCESS) { | ||
| 320 | + printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str()); | ||
| 321 | + return FAILED; | ||
| 322 | + } | ||
| 323 | + printf("%s - INFO - [XIR]: Finalize ir graph session success\n", GetTime().c_str()); | ||
| 324 | + return SUCCESS; | ||
| 325 | +} | ||
| @@ -0,0 +1,12 @@ | |||
| 1 | +# ---------------------------------------------------------------------------- | ||
| 2 | +# This program is free software, you can redistribute it and/or modify it. | ||
| 3 | +# Copyright (c) 2025 Huawei Technologies Co., Ltd. | ||
| 4 | +# This file is a part of the CANN Open Software. | ||
| 5 | +# Licensed under CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 6 | +# Please refer to the License for details. You may not use this file except in compliance with the License. | ||
| 7 | +# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING | ||
| 8 | +# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 9 | +# See LICENSE in the root of the software repository for the full text of the License. | ||
| 10 | +# ---------------------------------------------------------------------------- | ||
| 11 | + | ||
| 12 | +add_graph_plugin_sources() | ||
The file is empty
| @@ -0,0 +1,10 @@ | |||
| 1 | +# This program is free software, you can redistribute it and/or modify it. | ||
| 2 | +# Copyright (c) 2025 Huawei Technologies Co., Ltd. | ||
| 3 | +# This file is a part of the CANN Open Software. | ||
| 4 | +# Licensed under 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, INCLUDING | ||
| 7 | +# 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 | + | ||
| @@ -0,0 +1,46 @@ | |||
| 1 | +/** | ||
| 2 | + * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT) | ||
| 3 | + * and is contributed to the CANN Open Software. | ||
| 4 | + * | ||
| 5 | + * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT). | ||
| 6 | + * All Rights Reserved. | ||
| 7 | + * | ||
| 8 | + * Authors (accounts): | ||
| 9 | + * - Li Wen <@liwenkkklll> | ||
| 10 | + * - Su Tonghua <@sutonghua> | ||
| 11 | + * | ||
| 12 | + * This program is free software: you can redistribute it and/or modify it. | ||
| 13 | + * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 14 | + * You may not use this file except in compliance with the License. | ||
| 15 | + * See the LICENSE file at the root of the repository for the full text of the License. | ||
| 16 | + * | ||
| 17 | + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED, | ||
| 18 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 19 | + */ | ||
| 20 | + | ||
| 21 | +/*! | ||
| 22 | + * \file square_v2_graph_infer.cpp | ||
| 23 | + * \brief square_v2 operater graph infer resource | ||
| 24 | + */ | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | +namespace ops { | ||
| 29 | +using namespace ge; | ||
| 30 | + | ||
| 31 | +static constexpr int64_t IDX_0 = 0; | ||
| 32 | + | ||
| 33 | +static ge::graphStatus InferDataTypeSquareV2(gert::InferDataTypeContext* context) | ||
| 34 | +{ | ||
| 35 | + OP_LOGD(context->GetNodeName(), "Begin to do InferDataTypeSquareV2"); | ||
| 36 | + | ||
| 37 | + ge::DataType xDtype = context->GetInputDataType(IDX_0); | ||
| 38 | + context->SetOutputDataType(IDX_0, xDtype); | ||
| 39 | + | ||
| 40 | + OP_LOGD(context->GetNodeName(), "End to do InferDataTypeSquareV2"); | ||
| 41 | + return GRAPH_SUCCESS; | ||
| 42 | +} | ||
| 43 | + | ||
| 44 | +IMPL_OP(SquareV2).InferDataType(InferDataTypeSquareV2); | ||
| 45 | + | ||
| 46 | +}; // namespace ops | ||
| @@ -0,0 +1,50 @@ | |||
| 1 | +/** | ||
| 2 | + * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT) | ||
| 3 | + * and is contributed to the CANN Open Software. | ||
| 4 | + * | ||
| 5 | + * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT). | ||
| 6 | + * All Rights Reserved. | ||
| 7 | + * | ||
| 8 | + * Authors (accounts): | ||
| 9 | + * - Li Wen <@liwenkkklll> | ||
| 10 | + * - Su Tonghua <@sutonghua> | ||
| 11 | + * | ||
| 12 | + * This program is free software: you can redistribute it and/or modify it. | ||
| 13 | + * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 14 | + * You may not use this file except in compliance with the License. | ||
| 15 | + * See the LICENSE file at the root of the repository for the full text of the License. | ||
| 16 | + * | ||
| 17 | + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED, | ||
| 18 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 19 | + */ | ||
| 20 | + | ||
| 21 | +/*! | ||
| 22 | + * \file square_v2_proto.h | ||
| 23 | + * \brief | ||
| 24 | +*/ | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | +namespace ge { | ||
| 32 | +/** | ||
| 33 | +*@brief Returns element-wise square of "x". | ||
| 34 | + | ||
| 35 | +*@par Inputs: | ||
| 36 | +x: A ND Tensor of type float16 or float32. \n | ||
| 37 | +*@par Outputs: | ||
| 38 | +*z: A ND Tensor. Has the same dtype as "x". | ||
| 39 | +*@par Third-party framework compatibility | ||
| 40 | +*Compatible with the TensorFlow operator Square. | ||
| 41 | +*/ | ||
| 42 | + | ||
| 43 | +REG_OP(SquareV2) | ||
| 44 | + .INPUT(x, TensorType({DT_FLOAT, DT_FLOAT16})) | ||
| 45 | + .OUTPUT(z, TensorType({DT_FLOAT, DT_FLOAT16})) | ||
| 46 | + .OP_END_FACTORY_REG(SquareV2) | ||
| 47 | + | ||
| 48 | +} // namespace ge | ||
| 49 | + | ||
| 50 | + | ||
| @@ -0,0 +1,67 @@ | |||
| 1 | +{ | ||
| 2 | + "op_type": "SquareV2", | ||
| 3 | + "op_list": [ | ||
| 4 | + { | ||
| 5 | + "bin_filename": "SquareV2_a1532827238e1555db7b997c7bce2928", | ||
| 6 | + "inputs": [ | ||
| 7 | + { | ||
| 8 | + "name": "x", | ||
| 9 | + "index": 0, | ||
| 10 | + "dtype": "float32", | ||
| 11 | + "format": "ND", | ||
| 12 | + "paramType": "required", | ||
| 13 | + "shape": [ | ||
| 14 | + -2 | ||
| 15 | + ], | ||
| 16 | + "format_match_mode": "FormatAgnostic" | ||
| 17 | + } | ||
| 18 | + ], | ||
| 19 | + "attrs": [ | ||
| 20 | + ], | ||
| 21 | + "outputs": [ | ||
| 22 | + { | ||
| 23 | + "name": "z", | ||
| 24 | + "index": 0, | ||
| 25 | + "dtype": "float32", | ||
| 26 | + "format": "ND", | ||
| 27 | + "paramType": "required", | ||
| 28 | + "shape": [ | ||
| 29 | + -2 | ||
| 30 | + ], | ||
| 31 | + "format_match_mode": "FormatAgnostic" | ||
| 32 | + } | ||
| 33 | + ] | ||
| 34 | + }, | ||
| 35 | + { | ||
| 36 | + "bin_filename": "SquareV2_11132827238e1555db7b997c7bce2928", | ||
| 37 | + "inputs": [ | ||
| 38 | + { | ||
| 39 | + "name": "x", | ||
| 40 | + "index": 0, | ||
| 41 | + "dtype": "float16", | ||
| 42 | + "format": "ND", | ||
| 43 | + "paramType": "required", | ||
| 44 | + "shape": [ | ||
| 45 | + -2 | ||
| 46 | + ], | ||
| 47 | + "format_match_mode": "FormatAgnostic" | ||
| 48 | + } | ||
| 49 | + ], | ||
| 50 | + "attrs": [ | ||
| 51 | + ], | ||
| 52 | + "outputs": [ | ||
| 53 | + { | ||
| 54 | + "name": "z", | ||
| 55 | + "index": 0, | ||
| 56 | + "dtype": "float16", | ||
| 57 | + "format": "ND", | ||
| 58 | + "paramType": "required", | ||
| 59 | + "shape": [ | ||
| 60 | + -2 | ||
| 61 | + ], | ||
| 62 | + "format_match_mode": "FormatAgnostic" | ||
| 63 | + } | ||
| 64 | + ] | ||
| 65 | + } | ||
| 66 | + ] | ||
| 67 | +} | ||
| @@ -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 | +[SquareV2] | ||
| 13 | +default=0 | ||
| @@ -0,0 +1,58 @@ | |||
| 1 | +/** | ||
| 2 | + * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT) | ||
| 3 | + * and is contributed to the CANN Open Software. | ||
| 4 | + * | ||
| 5 | + * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT). | ||
| 6 | + * All Rights Reserved. | ||
| 7 | + * | ||
| 8 | + * Authors (accounts): | ||
| 9 | + * - Li Wen <@liwenkkklll> | ||
| 10 | + * - Su Tonghua <@sutonghua> | ||
| 11 | + * | ||
| 12 | + * This program is free software: you can redistribute it and/or modify it. | ||
| 13 | + * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 14 | + * You may not use this file except in compliance with the License. | ||
| 15 | + * See the LICENSE file at the root of the repository for the full text of the License. | ||
| 16 | + * | ||
| 17 | + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED, | ||
| 18 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 19 | + */ | ||
| 20 | + | ||
| 21 | +/*! | ||
| 22 | + * \file square_v2.cpp | ||
| 23 | + * \brief | ||
| 24 | + */ | ||
| 25 | + | ||
| 26 | + | ||
| 27 | +namespace ops { | ||
| 28 | + class SquareV2 : public OpDef { | ||
| 29 | + public: | ||
| 30 | + explicit SquareV2(const char* name) : OpDef(name) | ||
| 31 | + { | ||
| 32 | + this->Input("x") // 输入x定义 | ||
| 33 | + .ParamType(REQUIRED) // 必选输入 | ||
| 34 | + .DataType({ge::DT_FLOAT,ge::DT_FLOAT16}) // 支持数据类型 | ||
| 35 | + .Format({ge::FORMAT_ND,ge::FORMAT_ND}) // 支持format格式 | ||
| 36 | + .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND}) // 未确定大小shape对应format格式 | ||
| 37 | + .AutoContiguous(); // 内存自动连续化 | ||
| 38 | + | ||
| 39 | + this->Output("z") | ||
| 40 | + .ParamType(REQUIRED) | ||
| 41 | + .DataType({ge::DT_FLOAT, ge::DT_FLOAT16}) | ||
| 42 | + .Format({ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 43 | + .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND}) | ||
| 44 | + .AutoContiguous(); | ||
| 45 | + | ||
| 46 | + OpAICoreConfig aicoreConfig; | ||
| 47 | + aicoreConfig.DynamicCompileStaticFlag(true) | ||
| 48 | + .DynamicFormatFlag(false) | ||
| 49 | + .DynamicRankSupportFlag(true) | ||
| 50 | + .DynamicShapeSupportFlag(true) | ||
| 51 | + .NeedCheckSupportFlag(false) | ||
| 52 | + .PrecisionReduceFlag(true) | ||
| 53 | + .ExtendCfgInfo("opFile.value", "square_v2"); // 这里制定的值会对应到kernel入口文件名.cpp | ||
| 54 | + this->AICore().AddConfig("ascend910b", aicoreConfig); // 其他的soc版本补充部分配置项 | ||
| 55 | + } | ||
| 56 | + }; | ||
| 57 | + OP_ADD(SquareV2); // 添加算子信息库 | ||
| 58 | + } // namespace ops | ||
| @@ -0,0 +1,52 @@ | |||
| 1 | +/** | ||
| 2 | + * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT) | ||
| 3 | + * and is contributed to the CANN Open Software. | ||
| 4 | + * | ||
| 5 | + * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT). | ||
| 6 | + * All Rights Reserved. | ||
| 7 | + * | ||
| 8 | + * Authors (accounts): | ||
| 9 | + * - Li Wen <@liwenkkklll> | ||
| 10 | + * - Su Tonghua <@sutonghua> | ||
| 11 | + * | ||
| 12 | + * This program is free software: you can redistribute it and/or modify it. | ||
| 13 | + * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 14 | + * You may not use this file except in compliance with the License. | ||
| 15 | + * See the LICENSE file at the root of the repository for the full text of the License. | ||
| 16 | + * | ||
| 17 | + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED, | ||
| 18 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 19 | + */ | ||
| 20 | + | ||
| 21 | +/*! | ||
| 22 | + * \file square_v2_infer.cpp | ||
| 23 | + * \brief | ||
| 24 | + */ | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | +using namespace ge; | ||
| 29 | + | ||
| 30 | +namespace ops { | ||
| 31 | + static constexpr int64_t IDX_0 = 0; | ||
| 32 | + | ||
| 33 | + static ge::graphStatus InferShapeSquareV2(gert::InferShapeContext* context) | ||
| 34 | + { | ||
| 35 | + OP_LOGD(context->GetNodeName(), "Begin to do InferShapeSquareV2"); | ||
| 36 | + | ||
| 37 | + // get input shapes | ||
| 38 | + const gert::Shape* xShape = context->GetInputShape(IDX_0); | ||
| 39 | + OP_CHECK_NULL_WITH_CONTEXT(context, xShape); | ||
| 40 | + | ||
| 41 | + // get output shapes | ||
| 42 | + gert::Shape* zShape = context->GetOutputShape(IDX_0); | ||
| 43 | + OP_CHECK_NULL_WITH_CONTEXT(context, zShape); | ||
| 44 | + | ||
| 45 | + *zShape = *xShape; | ||
| 46 | + | ||
| 47 | + OP_LOGD(context->GetNodeName(), "End to do InferShapeSquareV2"); | ||
| 48 | + return GRAPH_SUCCESS; | ||
| 49 | + } | ||
| 50 | + | ||
| 51 | + IMPL_OP_INFERSHAPE(SquareV2).InferShape(InferShapeSquareV2); | ||
| 52 | + } // namespace ops | ||
| @@ -0,0 +1,227 @@ | |||
| 1 | +/** | ||
| 2 | + * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT) | ||
| 3 | + * and is contributed to the CANN Open Software. | ||
| 4 | + * | ||
| 5 | + * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT). | ||
| 6 | + * All Rights Reserved. | ||
| 7 | + * | ||
| 8 | + * Authors (accounts): | ||
| 9 | + * - Li Wen <@liwenkkklll> | ||
| 10 | + * - Su Tonghua <@sutonghua> | ||
| 11 | + * | ||
| 12 | + * This program is free software: you can redistribute it and/or modify it. | ||
| 13 | + * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 14 | + * You may not use this file except in compliance with the License. | ||
| 15 | + * See the LICENSE file at the root of the repository for the full text of the License. | ||
| 16 | + * | ||
| 17 | + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED, | ||
| 18 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 19 | + */ | ||
| 20 | + | ||
| 21 | +/*! | ||
| 22 | + * \file square_v2_tiling.cpp | ||
| 23 | + * \brief | ||
| 24 | + */ | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | +namespace optiling { | ||
| 34 | + | ||
| 35 | + using namespace Ops::Math::OpTiling; | ||
| 36 | + const uint32_t BLOCK_SIZE = 32; | ||
| 37 | + const uint32_t BUFFER_NUM = 2; | ||
| 38 | + const uint32_t WS_SYS_SIZE = 16U * 1024U * 1024U; | ||
| 39 | + | ||
| 40 | + struct SquareV2CompileInfo {}; | ||
| 41 | + | ||
| 42 | + // 获取平台信息如ubSize, coreNum | ||
| 43 | + static ge::graphStatus GetPlatformInfo(gert::TilingContext* context, uint64_t& ubSize, int64_t& coreNum) | ||
| 44 | + { | ||
| 45 | + // 获取ubsize coreNum | ||
| 46 | + fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo(); | ||
| 47 | + OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr); | ||
| 48 | + auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr); | ||
| 49 | + coreNum = ascendcPlatform.GetCoreNumAiv(); | ||
| 50 | + OP_CHECK_IF(coreNum == 0, OP_LOGE(context, "coreNum is 0"), return ge::GRAPH_FAILED); | ||
| 51 | + ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize); | ||
| 52 | + OP_CHECK_IF(ubSize == 0, OP_LOGE(context, "ubSize is 0"), return ge::GRAPH_FAILED); | ||
| 53 | + return ge::GRAPH_SUCCESS; | ||
| 54 | + } | ||
| 55 | + | ||
| 56 | + // 获取属性,shape信息 | ||
| 57 | + static ge::graphStatus GetShapeAttrsInfo(gert::TilingContext* context, int64_t& totalIdx, ge::DataType& dataType) | ||
| 58 | + { | ||
| 59 | + // 获取输入shape信息 | ||
| 60 | + auto inputX = context->GetInputShape(0); | ||
| 61 | + OP_CHECK_NULL_WITH_CONTEXT(context, inputX); | ||
| 62 | + // 如果输入shape 是标量 转换为{1},否则保持原 shape 不变 | ||
| 63 | + auto inputShapeX = EnsureNotScalar(inputX->GetStorageShape()); | ||
| 64 | + | ||
| 65 | + auto outZ = context->GetOutputShape(0); | ||
| 66 | + OP_CHECK_NULL_WITH_CONTEXT(context, outZ); | ||
| 67 | + auto outShapeZ = EnsureNotScalar(outZ->GetStorageShape()); | ||
| 68 | + | ||
| 69 | + // shape校验 | ||
| 70 | + bool shapeMatch = true; | ||
| 71 | + // 校验维度数一致 | ||
| 72 | + if (inputShapeX.GetDimNum() != outShapeZ.GetDimNum()) { | ||
| 73 | + shapeMatch = false; | ||
| 74 | + } else { | ||
| 75 | + // 校验每个维度的大小一致 | ||
| 76 | + size_t dimNum = inputShapeX.GetDimNum(); | ||
| 77 | + for (size_t i = 0; i < dimNum; i++) { | ||
| 78 | + if (inputShapeX.GetDim(i) != outShapeZ.GetDim(i)) { | ||
| 79 | + shapeMatch = false; | ||
| 80 | + break; | ||
| 81 | + } | ||
| 82 | + } | ||
| 83 | + } | ||
| 84 | + | ||
| 85 | + // 形状不匹配则报错 | ||
| 86 | + OP_CHECK_IF( | ||
| 87 | + !shapeMatch, | ||
| 88 | + OP_LOGE( | ||
| 89 | + context, "SquareV2: inputx,outputz shape not match! dim num: x=%zu, z=%zu", | ||
| 90 | + inputShapeX.GetDimNum(), outShapeZ.GetDimNum()), | ||
| 91 | + return ge::GRAPH_FAILED); | ||
| 92 | + | ||
| 93 | + totalIdx = inputX->GetOriginShape().GetShapeSize(); | ||
| 94 | + // dtype校验 | ||
| 95 | + const std::set<ge::DataType> supportedDtype = {ge::DT_FLOAT, ge::DT_FLOAT16}; | ||
| 96 | + auto inputDesc = context->GetInputDesc(0); | ||
| 97 | + OP_CHECK_NULL_WITH_CONTEXT(context, inputDesc); | ||
| 98 | + dataType = inputDesc->GetDataType(); | ||
| 99 | + if (supportedDtype.count(dataType) == 0) { | ||
| 100 | + OP_LOGE(context, "invalid dtype"); | ||
| 101 | + return ge::GRAPH_FAILED; | ||
| 102 | + } | ||
| 103 | + return ge::GRAPH_SUCCESS; | ||
| 104 | + } | ||
| 105 | + | ||
| 106 | + static ge::graphStatus GetWorkspaceSize(gert::TilingContext* context) | ||
| 107 | + { | ||
| 108 | + size_t* currentWorkspace = context->GetWorkspaceSizes(1); | ||
| 109 | + OP_CHECK_NULL_WITH_CONTEXT(context, currentWorkspace); | ||
| 110 | + currentWorkspace[0] = WS_SYS_SIZE; | ||
| 111 | + return ge::GRAPH_SUCCESS; | ||
| 112 | + } | ||
| 113 | + | ||
| 114 | + // tiling 分发入口 | ||
| 115 | + // 可直接替换你的 SquareV2TilingFunc 内部实现(保留函数签名) | ||
| 116 | + static ge::graphStatus SquareV2TilingFunc(gert::TilingContext* context) | ||
| 117 | + { | ||
| 118 | + // 1. platform | ||
| 119 | + uint64_t ubSize = 0; | ||
| 120 | + int64_t coreNum = 0; | ||
| 121 | + OP_CHECK_IF(GetPlatformInfo(context, ubSize, coreNum) != ge::GRAPH_SUCCESS, | ||
| 122 | + OP_LOGE(context, "GetPlatformInfo error"), return ge::GRAPH_FAILED); | ||
| 123 | + | ||
| 124 | + // 2. shapes & dtype | ||
| 125 | + int64_t totalIdx = 0; | ||
| 126 | + ge::DataType dataType; | ||
| 127 | + OP_CHECK_IF(GetShapeAttrsInfo(context, totalIdx, dataType) != ge::GRAPH_SUCCESS, | ||
| 128 | + OP_LOGE(context, "GetShapeAttrsInfo error"), return ge::GRAPH_FAILED); | ||
| 129 | + | ||
| 130 | + // 3. workspace | ||
| 131 | + OP_CHECK_IF(GetWorkspaceSize(context) != ge::GRAPH_SUCCESS, | ||
| 132 | + OP_LOGE(context, "GetWorkspaceSize error"), return ge::GRAPH_FAILED); | ||
| 133 | + | ||
| 134 | + SquareV2TilingData* tiling = context->GetTilingData<SquareV2TilingData>(); | ||
| 135 | + OP_CHECK_NULL_WITH_CONTEXT(context, tiling); | ||
| 136 | + OP_CHECK_IF(memset_s(tiling, sizeof(SquareV2TilingData), 0, sizeof(SquareV2TilingData)) != EOK, | ||
| 137 | + OP_LOGE(context, "set tiling data error"), return ge::GRAPH_FAILED); | ||
| 138 | + | ||
| 139 | + // --- safer numeric types --- | ||
| 140 | + uint32_t typeLength = 0; | ||
| 141 | + ge::TypeUtils::GetDataTypeLength(context->GetInputDesc(0)->GetDataType(), typeLength); | ||
| 142 | + if (typeLength == 0) { | ||
| 143 | + OP_LOGE(context, "typeLength is 0"); | ||
| 144 | + return ge::GRAPH_FAILED; | ||
| 145 | + } | ||
| 146 | + uint64_t inputBytes = static_cast<uint64_t>(typeLength); | ||
| 147 | + uint64_t inputLengthBytes = static_cast<uint64_t>(totalIdx) * inputBytes; | ||
| 148 | + | ||
| 149 | + // ub-based tileBlockNum guard (避免为0) | ||
| 150 | + uint32_t ubDataNumber = (inputBytes == 1ULL) ? 4U : 2U; | ||
| 151 | + uint64_t tmp = (ubSize / BLOCK_SIZE / BUFFER_NUM); | ||
| 152 | + uint32_t tileBlockNum = 1U; | ||
| 153 | + if (tmp > 0) { | ||
| 154 | + uint64_t tb = tmp / ubDataNumber; | ||
| 155 | + tileBlockNum = (tb == 0) ? 1U : static_cast<uint32_t>(tb); | ||
| 156 | + } | ||
| 157 | + | ||
| 158 | + // 每个 tile 包含的元素数(至少 1) | ||
| 159 | + uint32_t tileDataNum = static_cast<uint32_t>((static_cast<uint64_t>(tileBlockNum) * BLOCK_SIZE) / inputBytes); | ||
| 160 | + if (tileDataNum == 0U) tileDataNum = 1U; | ||
| 161 | + | ||
| 162 | + // 总 block 数(向上取整) | ||
| 163 | + uint64_t blocksTotal = (inputLengthBytes + BLOCK_SIZE - 1ULL) / BLOCK_SIZE; | ||
| 164 | + uint64_t coreNum64 = static_cast<uint64_t>(coreNum); | ||
| 165 | + if (coreNum64 > blocksTotal){ | ||
| 166 | + coreNum64 = blocksTotal; | ||
| 167 | + } | ||
| 168 | + if (coreNum64 == 0ULL) coreNum64 = 1ULL; // 最少 1 core | ||
| 169 | + uint32_t finalCoreNum = static_cast<uint32_t>(coreNum64); | ||
| 170 | + | ||
| 171 | + uint64_t everyCoreInputBlockNum = blocksTotal / coreNum64; // 基本块数 | ||
| 172 | + uint32_t tailBlockNum = static_cast<uint32_t>(blocksTotal % coreNum64); // 前 tailBlockNum 个核是 big-core | ||
| 173 | + | ||
| 174 | + // small-core 数量(元素) | ||
| 175 | + uint64_t smallCoreDataNum_u = everyCoreInputBlockNum * BLOCK_SIZE / inputBytes; | ||
| 176 | + uint32_t smallCoreDataNum = static_cast<uint32_t>(smallCoreDataNum_u); | ||
| 177 | + | ||
| 178 | + uint32_t smallTileNum = static_cast<uint32_t>(everyCoreInputBlockNum / static_cast<uint64_t>(tileBlockNum)); | ||
| 179 | + uint32_t finalSmallTileNum = ((everyCoreInputBlockNum % tileBlockNum) == 0) ? smallTileNum : (smallTileNum + 1); | ||
| 180 | + int64_t smallTailDataNum_i = static_cast<int64_t>(smallCoreDataNum) - static_cast<int64_t>(tileDataNum) * static_cast<int64_t>(smallTileNum); | ||
| 181 | + uint32_t smallTailDataNum = (smallTailDataNum_i <= 0) ? tileDataNum : static_cast<uint32_t>(smallTailDataNum_i); | ||
| 182 | + | ||
| 183 | + // big-core(每个多一个 block) | ||
| 184 | + uint64_t bigEveryCoreBlockNum = everyCoreInputBlockNum + 1ULL; | ||
| 185 | + uint64_t bigCoreDataNum_u = bigEveryCoreBlockNum * BLOCK_SIZE / inputBytes; | ||
| 186 | + uint32_t bigCoreDataNum = static_cast<uint32_t>(bigCoreDataNum_u); | ||
| 187 | + uint32_t bigTileNum = static_cast<uint32_t>(bigEveryCoreBlockNum / static_cast<uint64_t>(tileBlockNum)); | ||
| 188 | + uint32_t finalBigTileNum = ((bigEveryCoreBlockNum % tileBlockNum) == 0) ? bigTileNum : (bigTileNum + 1); | ||
| 189 | + int64_t bigTailDataNum_i = static_cast<int64_t>(bigCoreDataNum) - static_cast<int64_t>(tileDataNum) * static_cast<int64_t>(bigTileNum); | ||
| 190 | + uint32_t bigTailDataNum = (bigTailDataNum_i <= 0) ? tileDataNum : static_cast<uint32_t>(bigTailDataNum_i); | ||
| 191 | + | ||
| 192 | + // write back | ||
| 193 | + tiling->smallCoreDataNum = static_cast<int64_t>(smallCoreDataNum); | ||
| 194 | + tiling->bigCoreDataNum = static_cast<int64_t>(bigCoreDataNum); | ||
| 195 | + tiling->tileDataNum = static_cast<int64_t>(tileDataNum); | ||
| 196 | + tiling->smallTailDataNum = static_cast<int64_t>(smallTailDataNum); | ||
| 197 | + tiling->bigTailDataNum = static_cast<int64_t>(bigTailDataNum); | ||
| 198 | + tiling->finalSmallTileNum = static_cast<int64_t>(finalSmallTileNum); | ||
| 199 | + tiling->finalBigTileNum = static_cast<int64_t>(finalBigTileNum); | ||
| 200 | + tiling->tailBlockNum = static_cast<int64_t>(tailBlockNum); | ||
| 201 | + | ||
| 202 | + context->SetBlockDim(finalCoreNum); | ||
| 203 | + | ||
| 204 | + uint64_t tilingKey = 0; | ||
| 205 | + if (dataType == ge::DT_FLOAT) { | ||
| 206 | + tilingKey = GET_TPL_TILING_KEY(ELEMENTWISE_TPL_SCH_MODE_0); | ||
| 207 | + context->SetTilingKey(tilingKey); | ||
| 208 | + } else if (dataType == ge::DT_FLOAT16) { | ||
| 209 | + tilingKey = GET_TPL_TILING_KEY(ELEMENTWISE_TPL_SCH_MODE_1); | ||
| 210 | + context->SetTilingKey(tilingKey); | ||
| 211 | + } else { | ||
| 212 | + OP_LOGE(context, "get dtype error"); | ||
| 213 | + return ge::GRAPH_FAILED; | ||
| 214 | + } | ||
| 215 | + return ge::GRAPH_SUCCESS; | ||
| 216 | + } | ||
| 217 | + | ||
| 218 | + | ||
| 219 | + static ge::graphStatus TilingParseForSquareV2([[maybe_unused]] gert::TilingParseContext* context) | ||
| 220 | + { | ||
| 221 | + return ge::GRAPH_SUCCESS; | ||
| 222 | + } | ||
| 223 | + | ||
| 224 | + // tiling注册入口. | ||
| 225 | + IMPL_OP_OPTILING(SquareV2).Tiling(SquareV2TilingFunc).TilingParse<SquareV2CompileInfo>(TilingParseForSquareV2); | ||
| 226 | + } // namespace optiling | ||
| 227 | + | ||
| @@ -0,0 +1,50 @@ | |||
| 1 | +/** | ||
| 2 | + * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT) | ||
| 3 | + * and is contributed to the CANN Open Software. | ||
| 4 | + * | ||
| 5 | + * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT). | ||
| 6 | + * All Rights Reserved. | ||
| 7 | + * | ||
| 8 | + * Authors (accounts): | ||
| 9 | + * - Li Wen <@liwenkkklll> | ||
| 10 | + * - Su Tonghua <@sutonghua> | ||
| 11 | + * | ||
| 12 | + * This program is free software: you can redistribute it and/or modify it. | ||
| 13 | + * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 14 | + * You may not use this file except in compliance with the License. | ||
| 15 | + * See the LICENSE file at the root of the repository for the full text of the License. | ||
| 16 | + * | ||
| 17 | + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED, | ||
| 18 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 19 | + */ | ||
| 20 | + | ||
| 21 | +/*! | ||
| 22 | + * \file square_v2.cpp | ||
| 23 | + * \brief | ||
| 24 | + */ | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | +enum class SquareV2TilingKey : uint32_t | ||
| 29 | +{ | ||
| 30 | + TILING_KEY_EXAMPLE_FLOAT = 0, | ||
| 31 | + TILING_KEY_EXAMPLE_HALF = 1, | ||
| 32 | +}; | ||
| 33 | + | ||
| 34 | +template <uint32_t schMode> | ||
| 35 | +__global__ __aicore__ void square_v2(GM_ADDR x, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) | ||
| 36 | +{ | ||
| 37 | + REGISTER_TILING_DEFAULT(SquareV2TilingData); | ||
| 38 | + GET_TILING_DATA_WITH_STRUCT(SquareV2TilingData, tilingData, tiling); | ||
| 39 | + if constexpr (schMode == static_cast<uint32_t>(SquareV2TilingKey::TILING_KEY_EXAMPLE_FLOAT)) { | ||
| 40 | + NsSquareV2::SquareV2<float> op; // 算子kernel实例获取 | ||
| 41 | + op.Init(x, z, &tilingData); // 算子kernel实例初始化 | ||
| 42 | + op.Process(); // 算子kernel实例执行 | ||
| 43 | + } | ||
| 44 | + | ||
| 45 | + else if constexpr (schMode == static_cast<uint32_t>(SquareV2TilingKey::TILING_KEY_EXAMPLE_HALF)) { | ||
| 46 | + NsSquareV2::SquareV2<half> op; // 算子kernel实例获取 | ||
| 47 | + op.Init(x, z, &tilingData); // 算子kernel实例初始化 | ||
| 48 | + op.Process(); // 算子kernel实例执行 | ||
| 49 | + } | ||
| 50 | +} | ||
| @@ -0,0 +1,132 @@ | |||
| 1 | +/** | ||
| 2 | + * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT) | ||
| 3 | + * and is contributed to the CANN Open Software. | ||
| 4 | + * | ||
| 5 | + * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT). | ||
| 6 | + * All Rights Reserved. | ||
| 7 | + * | ||
| 8 | + * Authors (accounts): | ||
| 9 | + * - Li Wen <@liwenkkklll> | ||
| 10 | + * - Su Tonghua <@sutonghua> | ||
| 11 | + * | ||
| 12 | + * This program is free software: you can redistribute it and/or modify it. | ||
| 13 | + * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 14 | + * You may not use this file except in compliance with the License. | ||
| 15 | + * See the LICENSE file at the root of the repository for the full text of the License. | ||
| 16 | + * | ||
| 17 | + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED, | ||
| 18 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 19 | + */ | ||
| 20 | + | ||
| 21 | +/*! | ||
| 22 | + * \file square_v2.h | ||
| 23 | + * \brief | ||
| 24 | + */ | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | + | ||
| 32 | + | ||
| 33 | +namespace NsSquareV2 { | ||
| 34 | + | ||
| 35 | + using namespace AscendC; | ||
| 36 | + | ||
| 37 | + constexpr int32_t BUFFER_NUM = 2; | ||
| 38 | + | ||
| 39 | + template <typename T> | ||
| 40 | + class SquareV2 { | ||
| 41 | + public: | ||
| 42 | + __aicore__ inline SquareV2(){}; | ||
| 43 | + | ||
| 44 | + __aicore__ inline void Init(GM_ADDR x, GM_ADDR z, const SquareV2TilingData* tilingData); | ||
| 45 | + __aicore__ inline void Process(); | ||
| 46 | + | ||
| 47 | + private: | ||
| 48 | + __aicore__ inline void CopyIn(int32_t progress); | ||
| 49 | + __aicore__ inline void CopyOut(int32_t progress); | ||
| 50 | + __aicore__ inline void Compute(int32_t progress); | ||
| 51 | + | ||
| 52 | + private: | ||
| 53 | + TPipe pipe; | ||
| 54 | + TQue<QuePosition::VECIN, BUFFER_NUM> inputQueueX; | ||
| 55 | + TQue<QuePosition::VECOUT, BUFFER_NUM> outputQueueZ; | ||
| 56 | + GlobalTensor<T> inputGMX; | ||
| 57 | + GlobalTensor<T> outputGMZ; | ||
| 58 | + | ||
| 59 | + uint32_t coreDataNum; | ||
| 60 | + uint32_t tileNum; | ||
| 61 | + uint32_t tileDataNum; | ||
| 62 | + uint32_t tailDataNum; | ||
| 63 | + uint32_t processDataNum; | ||
| 64 | + }; | ||
| 65 | + | ||
| 66 | + template <typename T> | ||
| 67 | + __aicore__ inline void SquareV2<T>::Init(GM_ADDR x, GM_ADDR z, const SquareV2TilingData* tilingData) | ||
| 68 | + { | ||
| 69 | + ASSERT(AscendC::GetBlockNum() != 0 && "block dim can not be zero!"); | ||
| 70 | + uint32_t coreNum = AscendC::GetBlockIdx(); | ||
| 71 | + uint32_t globalBufferIndex = tilingData->bigCoreDataNum * AscendC::GetBlockIdx(); | ||
| 72 | + this->tileDataNum = tilingData->tileDataNum; | ||
| 73 | + if (coreNum < tilingData->tailBlockNum) { | ||
| 74 | + this->coreDataNum = tilingData->bigCoreDataNum; | ||
| 75 | + this->tileNum = tilingData->finalBigTileNum; | ||
| 76 | + this->tailDataNum = tilingData->bigTailDataNum; | ||
| 77 | + } | ||
| 78 | + else { | ||
| 79 | + this->coreDataNum = tilingData->smallCoreDataNum; | ||
| 80 | + this->tileNum = tilingData->finalSmallTileNum; | ||
| 81 | + this->tailDataNum = tilingData->smallTailDataNum; | ||
| 82 | + globalBufferIndex -= (tilingData->bigCoreDataNum - tilingData->smallCoreDataNum) * (AscendC::GetBlockIdx() - tilingData->tailBlockNum); | ||
| 83 | + } | ||
| 84 | + inputGMX.SetGlobalBuffer((__gm__ T*)x + globalBufferIndex, this->coreDataNum); | ||
| 85 | + outputGMZ.SetGlobalBuffer((__gm__ T*)z + globalBufferIndex, this->coreDataNum); | ||
| 86 | + pipe.InitBuffer(inputQueueX, BUFFER_NUM, this->tileDataNum * sizeof(T)); | ||
| 87 | + pipe.InitBuffer(outputQueueZ, BUFFER_NUM, this->tileDataNum * sizeof(T)); | ||
| 88 | + } | ||
| 89 | + | ||
| 90 | + template <typename T> | ||
| 91 | + __aicore__ inline void SquareV2<T>::CopyIn(int32_t progress) | ||
| 92 | + { | ||
| 93 | + AscendC::LocalTensor<T> xLocal = inputQueueX.AllocTensor<T>(); | ||
| 94 | + AscendC::DataCopy(xLocal, inputGMX[progress * this->tileDataNum], this->processDataNum); | ||
| 95 | + inputQueueX.EnQue(xLocal); | ||
| 96 | + } | ||
| 97 | + | ||
| 98 | + template <typename T> | ||
| 99 | + __aicore__ inline void SquareV2<T>::CopyOut(int32_t progress) | ||
| 100 | + { | ||
| 101 | + AscendC::LocalTensor<T> zLocal = outputQueueZ.DeQue<T>(); | ||
| 102 | + AscendC::DataCopy(outputGMZ[progress * this->tileDataNum], zLocal, this->processDataNum); | ||
| 103 | + outputQueueZ.FreeTensor(zLocal); | ||
| 104 | + } | ||
| 105 | + | ||
| 106 | + template <typename T> | ||
| 107 | + __aicore__ inline void SquareV2<T>::Compute(int32_t progress) | ||
| 108 | + { | ||
| 109 | + AscendC::LocalTensor<T> xLocal = inputQueueX.DeQue<T>(); | ||
| 110 | + AscendC::LocalTensor<T> zLocal = outputQueueZ.AllocTensor<T>(); | ||
| 111 | + AscendC::Mul(zLocal, xLocal,xLocal, this->processDataNum); | ||
| 112 | + outputQueueZ.EnQue<T>(zLocal); | ||
| 113 | + inputQueueX.FreeTensor(xLocal); | ||
| 114 | + } | ||
| 115 | + | ||
| 116 | + template <typename T> | ||
| 117 | + __aicore__ inline void SquareV2<T>::Process() | ||
| 118 | + { | ||
| 119 | + int32_t loopCount = this->tileNum; | ||
| 120 | + this->processDataNum = this->tileDataNum; | ||
| 121 | + for (int32_t i = 0; i < loopCount; i++) { | ||
| 122 | + if (i == this->tileNum - 1) { | ||
| 123 | + this->processDataNum = this->tailDataNum; | ||
| 124 | + } | ||
| 125 | + CopyIn(i); | ||
| 126 | + Compute(i); | ||
| 127 | + CopyOut(i); | ||
| 128 | + } | ||
| 129 | + } | ||
| 130 | + | ||
| 131 | +} // namespace NsSquareV2 | ||
| 132 | + | ||
| @@ -0,0 +1,39 @@ | |||
| 1 | +/** | ||
| 2 | + * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT) | ||
| 3 | + * and is contributed to the CANN Open Software. | ||
| 4 | + * | ||
| 5 | + * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT). | ||
| 6 | + * All Rights Reserved. | ||
| 7 | + * | ||
| 8 | + * Authors (accounts): | ||
| 9 | + * - Li Wen <@liwenkkklll> | ||
| 10 | + * - Su Tonghua <@sutonghua> | ||
| 11 | + * | ||
| 12 | + * This program is free software: you can redistribute it and/or modify it. | ||
| 13 | + * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 14 | + * You may not use this file except in compliance with the License. | ||
| 15 | + * See the LICENSE file at the root of the repository for the full text of the License. | ||
| 16 | + * | ||
| 17 | + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED, | ||
| 18 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 19 | + */ | ||
| 20 | + | ||
| 21 | +/*! | ||
| 22 | + * \file square_v2_tiling_data.h | ||
| 23 | + * \brief tiling data struct | ||
| 24 | + */ | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | + | ||
| 29 | +struct SquareV2TilingData { | ||
| 30 | + int64_t smallCoreDataNum; | ||
| 31 | + int64_t bigCoreDataNum; | ||
| 32 | + int64_t finalBigTileNum; | ||
| 33 | + int64_t finalSmallTileNum; | ||
| 34 | + int64_t tileDataNum; | ||
| 35 | + int64_t smallTailDataNum; | ||
| 36 | + int64_t bigTailDataNum; | ||
| 37 | + int64_t tailBlockNum; | ||
| 38 | +}; | ||
| 39 | + | ||
| @@ -0,0 +1,45 @@ | |||
| 1 | +/** | ||
| 2 | + * This file is part of the OpenBOAT project at Harbin Institute of Technology (HIT) | ||
| 3 | + * and is contributed to the CANN Open Software. | ||
| 4 | + * | ||
| 5 | + * Copyright (c) 2025 AISS Group, Harbin Institute of Technology (HIT). | ||
| 6 | + * All Rights Reserved. | ||
| 7 | + * | ||
| 8 | + * Authors (accounts): | ||
| 9 | + * - Li Wen <@liwenkkklll> | ||
| 10 | + * - Su Tonghua <@sutonghua> | ||
| 11 | + * | ||
| 12 | + * This program is free software: you can redistribute it and/or modify it. | ||
| 13 | + * Licensed under the CANN Open Software License Agreement Version 2.0 (the "License"). | ||
| 14 | + * You may not use this file except in compliance with the License. | ||
| 15 | + * See the LICENSE file at the root of the repository for the full text of the License. | ||
| 16 | + * | ||
| 17 | + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTIES OF ANY KIND, EXPRESS OR IMPLIED, | ||
| 18 | + * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. | ||
| 19 | + */ | ||
| 20 | + | ||
| 21 | +/*! | ||
| 22 | + * \file square_v2_tiling_key.h | ||
| 23 | + * \brief square_v2 tiling key declare | ||
| 24 | + */ | ||
| 25 | + | ||
| 26 | + | ||
| 27 | + | ||
| 28 | + | ||
| 29 | + | ||
| 30 | + | ||
| 31 | +/* Mode场景定义 */ | ||
| 32 | + | ||
| 33 | + | ||
| 34 | +/* 继续定义其他Mode场景... */ | ||
| 35 | + | ||
| 36 | +/* 模板参数 */ | ||
| 37 | +ASCENDC_TPL_ARGS_DECL( | ||
| 38 | + SquareV2, | ||
| 39 | + ASCENDC_TPL_UINT_DECL(schMode, 1, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1)); | ||
| 40 | + | ||
| 41 | +/* 模板参数组合 */ | ||
| 42 | +ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL( | ||
| 43 | + ASCENDC_TPL_UINT_SEL(schMode, ASCENDC_TPL_UI_LIST, ELEMENTWISE_TPL_SCH_MODE_0, ELEMENTWISE_TPL_SCH_MODE_1))); | ||
| 44 | + | ||
| 45 | + | ||
The file is empty