已合并
fix: 修复opgen生成的算子工程无法编译通过的问题 #2431
songkai111创建于 4月24日
fix: 修复opgen生成的算子工程无法编译通过的问题 #2431
已合并
共 6 个文件变更+52-13
| @@ -118,10 +118,16 @@ endif() | |||
| 118 | 118 | ||
| 119 | set(_orig_ops "${ASCEND_OP_NAME}") | 119 | set(_orig_ops "${ASCEND_OP_NAME}") |
| 120 | list(LENGTH _orig_ops _orig_len) | 120 | list(LENGTH _orig_ops _orig_len) |
| 121 | -set(_filtered_ops "${_orig_ops}") | 121 | +set(_include_examples FALSE) |
| 122 | -list(FILTER _filtered_ops INCLUDE REGEX "add_example") | 122 | +if(_orig_len GREATER 0) |
| 123 | -list(LENGTH _filtered_ops _filtered_len) | 123 | + foreach(_op ${_orig_ops}) |
| 124 | -if(_orig_len GREATER 0 AND _filtered_len EQUAL _orig_len) | 124 | + if(EXISTS ${CMAKE_CURRENT_SOURCE_DIR}/examples/${_op}/CMakeLists.txt) |
| 125 | + set(_include_examples TRUE) | ||
| 126 | + break() | ||
| 127 | + endif() | ||
| 128 | + endforeach() | ||
| 129 | +endif() | ||
| 130 | +if(_include_examples) | ||
| 125 | add_subdirectory(examples) | 131 | add_subdirectory(examples) |
| 126 | endif() | 132 | endif() |
| 127 | 133 | ||
| @@ -37,6 +37,7 @@ class OpGenerator: | |||
| 37 | self._copy_template() | 37 | self._copy_template() |
| 38 | self._rename_files() | 38 | self._rename_files() |
| 39 | self._replace_content() | 39 | self._replace_content() |
| 40 | + self._update_cmake_chain() | ||
| 40 | logging.info(f"成功为 {self.op_type}/{self.op_name} 创建算子工程!") | 41 | logging.info(f"成功为 {self.op_type}/{self.op_name} 创建算子工程!") |
| 41 | logging.info(f"工程路径: {self.dest_dir}") | 42 | logging.info(f"工程路径: {self.dest_dir}") |
| 42 | 43 | ||
| @@ -80,6 +81,36 @@ class OpGenerator: | |||
| 80 | except OSError as e: | 81 | except OSError as e: |
| 81 | raise OSError(f"重命名 '{old_path}' 到 '{new_path}' 失败: {e}") from e | 82 | raise OSError(f"重命名 '{old_path}' 到 '{new_path}' 失败: {e}") from e |
| 82 | 83 | ||
| 84 | + def _update_cmake_chain(self): | ||
| 85 | + """在父目录的CMakeLists.txt中添加add_subdirectory,确保构建系统能找到新目录""" | ||
| 86 | + current = self.dest_dir | ||
| 87 | + project_root = os.path.abspath(self.output_path) | ||
| 88 | + | ||
| 89 | + while True: | ||
| 90 | + parent = os.path.dirname(current) | ||
| 91 | + child = os.path.basename(current) | ||
| 92 | + parent_abs = os.path.abspath(parent) | ||
| 93 | + | ||
| 94 | + cmake_file = os.path.join(parent, "CMakeLists.txt") | ||
| 95 | + if os.path.exists(cmake_file): | ||
| 96 | + with open(cmake_file, 'r', encoding='utf-8') as f: | ||
| 97 | + content = f.read() | ||
| 98 | + | ||
| 99 | + already_included = bool(re.search( | ||
| 100 | + rf'add_subdirectory\s*\(\s*{re.escape(child)}\s*[\s\)]', content | ||
| 101 | + )) | ||
| 102 | + uses_glob = bool(re.search(r'file\s*\(\s*GLOB', content)) | ||
| 103 | + | ||
| 104 | + if not already_included and not uses_glob: | ||
| 105 | + with open(cmake_file, 'a', encoding='utf-8') as f: | ||
| 106 | + f.write(f"\nadd_subdirectory({child})\n") | ||
| 107 | + logging.info(f"Added add_subdirectory({child}) to {cmake_file}") | ||
| 108 | + | ||
| 109 | + if parent_abs == project_root: | ||
| 110 | + break | ||
| 111 | + | ||
| 112 | + current = parent | ||
| 113 | + | ||
| 83 | def _replace_content_in_file(self, file_path, replacements): | 114 | def _replace_content_in_file(self, file_path, replacements): |
| 84 | """Helper to replace content in a single file.""" | 115 | """Helper to replace content in a single file.""" |
| 85 | try: | 116 | try: |
| @@ -24,11 +24,11 @@ static ge::graphStatus InferShapeAddExample(gert::InferShapeContext* context) | |||
| 24 | OP_LOGD(context->GetNodeName(), "Begin to do InferShapeAddExample"); | 24 | OP_LOGD(context->GetNodeName(), "Begin to do InferShapeAddExample"); |
| 25 | 25 | ||
| 26 | // get input shapes | 26 | // get input shapes |
| 27 | - const gert::Shape* xShape = context->GetInputShape(IDX_0); | 27 | + const gert::Shape* xShape = context->GetInputShape(0); |
| 28 | OP_CHECK_NULL_WITH_CONTEXT(context, xShape); | 28 | OP_CHECK_NULL_WITH_CONTEXT(context, xShape); |
| 29 | 29 | ||
| 30 | // get output shapes | 30 | // get output shapes |
| 31 | - gert::Shape* yShape = context->GetOutputShape(IDX_0); | 31 | + gert::Shape* yShape = context->GetOutputShape(0); |
| 32 | OP_CHECK_NULL_WITH_CONTEXT(context, yShape); | 32 | OP_CHECK_NULL_WITH_CONTEXT(context, yShape); |
| 33 | 33 | ||
| 34 | // 填充输出shape大小 | 34 | // 填充输出shape大小 |
| @@ -25,8 +25,9 @@ namespace optiling { | |||
| 25 | struct AddExampleCompileInfo {}; | 25 | struct AddExampleCompileInfo {}; |
| 26 | 26 | ||
| 27 | // tiling 分发入口 | 27 | // tiling 分发入口 |
| 28 | -static ge::graphStatus AddExampleTilingFunc(gert::TilingContext* context) | 28 | +static ge::graphStatus AddExampleTilingFunc([[maybe_unused]] gert::TilingContext* context) |
| 29 | { | 29 | { |
| 30 | + return ge::GRAPH_SUCCESS; | ||
| 30 | } | 31 | } |
| 31 | 32 | ||
| 32 | static ge::graphStatus TilingParseForAddExample([[maybe_unused]] gert::TilingParseContext* context) | 33 | static ge::graphStatus TilingParseForAddExample([[maybe_unused]] gert::TilingParseContext* context) |
| @@ -22,7 +22,7 @@ enum class AddExampleTilingKey : uint32_t | |||
| 22 | }; | 22 | }; |
| 23 | 23 | ||
| 24 | template <uint32_t schMode> | 24 | template <uint32_t schMode> |
| 25 | -__global__ __aicore__ void add_example(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR workspace, GM_ADDR tiling) | 25 | +__global__ __aicore__ void add_example(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) |
| 26 | { | 26 | { |
| 27 | REGISTER_TILING_DEFAULT(AddExampleTilingData); | 27 | REGISTER_TILING_DEFAULT(AddExampleTilingData); |
| 28 | GET_TILING_DATA_WITH_STRUCT(AddExampleTilingData, tilingData, tiling); | 28 | GET_TILING_DATA_WITH_STRUCT(AddExampleTilingData, tilingData, tiling); |
| @@ -30,14 +30,14 @@ __global__ __aicore__ void add_example(GM_ADDR x, GM_ADDR y, GM_ADDR z, GM_ADDR | |||
| 30 | // 场景1 | 30 | // 场景1 |
| 31 | if constexpr (schMode == static_cast<uint32_t>(AddExampleTilingKey::TILING_KEY_EXAMPLE_FLOAT)) { | 31 | if constexpr (schMode == static_cast<uint32_t>(AddExampleTilingKey::TILING_KEY_EXAMPLE_FLOAT)) { |
| 32 | NsAddExample::AddExample<float> op; // 算子kernel实例获取 | 32 | NsAddExample::AddExample<float> op; // 算子kernel实例获取 |
| 33 | - op.Init(x, y, z, &tilingData); // 算子kernel实例初始化 | 33 | + op.Init(x, y, &tilingData); // 算子kernel实例初始化 |
| 34 | op.Process(); // 算子kernel实例执行 | 34 | op.Process(); // 算子kernel实例执行 |
| 35 | } | 35 | } |
| 36 | 36 | ||
| 37 | // 场景2 | 37 | // 场景2 |
| 38 | if constexpr (schMode == static_cast<uint32_t>(AddExampleTilingKey::TILING_KEY_EXAMPLE_INT32)) { | 38 | if constexpr (schMode == static_cast<uint32_t>(AddExampleTilingKey::TILING_KEY_EXAMPLE_INT32)) { |
| 39 | NsAddExample::AddExample<int32_t> op; // 算子kernel实例获取 | 39 | NsAddExample::AddExample<int32_t> op; // 算子kernel实例获取 |
| 40 | - op.Init(x, y, z, &tilingData); // 算子kernel实例初始化 | 40 | + op.Init(x, y, &tilingData); // 算子kernel实例初始化 |
| 41 | op.Process(); // 算子kernel实例执行 | 41 | op.Process(); // 算子kernel实例执行 |
| 42 | } | 42 | } |
| 43 | } | 43 | } |
| @@ -31,7 +31,7 @@ class AddExample { | |||
| 31 | public: | 31 | public: |
| 32 | __aicore__ inline AddExample(){}; | 32 | __aicore__ inline AddExample(){}; |
| 33 | 33 | ||
| 34 | - __aicore__ inline void Init(/*参数列表*/); | 34 | + __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, const AddExampleTilingData* tilingData/*参数列表*/); |
| 35 | __aicore__ inline void Process(/*参数列表*/); | 35 | __aicore__ inline void Process(/*参数列表*/); |
| 36 | 36 | ||
| 37 | private: | 37 | private: |
| @@ -46,7 +46,7 @@ private: | |||
| 46 | }; | 46 | }; |
| 47 | 47 | ||
| 48 | template <typename T> | 48 | template <typename T> |
| 49 | -__aicore__ inline void AddExample<T>::Init(/*参数列表*/) | 49 | +__aicore__ inline void AddExample<T>::Init(GM_ADDR x, GM_ADDR y, const AddExampleTilingData* tilingData/*参数列表*/) |
| 50 | { | 50 | { |
| 51 | } | 51 | } |
| 52 | 52 | ||
| @@ -68,7 +68,8 @@ __aicore__ inline void AddExample<T>::Compute(/*参数列表*/) | |||
| 68 | template <typename T> | 68 | template <typename T> |
| 69 | __aicore__ inline void AddExample<T>::Process() | 69 | __aicore__ inline void AddExample<T>::Process() |
| 70 | { | 70 | { |
| 71 | - for (int32_t i = 0; i < /*循环次数*/; i++) { | 71 | + int32_t loopExample = 5; |
| 72 | + for (int32_t i = 0; i < loopExample/*循环次数*/; i++) { | ||
| 72 | CopyIn(/*参数列表*/); | 73 | CopyIn(/*参数列表*/); |
| 73 | Compute(/*参数列表*/); | 74 | Compute(/*参数列表*/); |
| 74 | CopyOut(/*参数列表*/); | 75 | CopyOut(/*参数列表*/); |