feat: Python 自定义算子原型执行样例-ge-AtomGit

已合并

feat: Python 自定义算子原型执行样例 #4519

lfz2812创建于 26 天前从lfz2812/ge： develop4-pr3-base 合入到 cann/ge： develop

feat: Python 自定义算子原型执行样例 #4519

已合并

lfz2812创建于 26 天前

从lfz2812/ge develop4-pr3-base合入到cann/ge develop

收起

docs
- en / design / modules / ge_python
  - ge_python_custom_op_design.md
- zh
  - api / graph_engine_api
    - README.md
    - python / ge / runtime
      - Tensor
        Tensor.md
        addr.md
        data_type.md
        expand_dims_type.md
        format.md
        origin_format.md
        origin_shape.md
        overview.md
        placement.md
        shape.md
        shape_size.md
        size.md
        storage_format.md
        storage_shape.md
      - TensorDesc
        TensorDesc.md
        TensorDesc_constructor.md
        data_type.md
        overview.md
        shape.md
      - overview.md
  - design / modules / ge_python
    - ge_python_custom_op_design.md
examples / custom_op
- README.md
- README_en.md
- annotated_args_refresh_add_custom
  - README.md
  - README_en.md
  - offline
    - cpp
      - CMakeLists.txt
      - README.md
      - README_en.md
      - ge
        add_custom_ir.h
        add_custom_kernel.cpp
        custom_op.cpp
        utils
        compile_utils.cpp
        kernel_binary_map_utils.cpp
      - graph_build
        main.cc
      - model_exec
        main.cc
      - run.sh
    - python
      - CMakeLists.txt
      - README.md
      - README_en.md
      - proto
        add_custom.cc
        add_custom.h
      - run.sh
      - src
        build_graph.py
        ge
        annotated_add_custom.py
        run_model.py
  - online
    - cpp
      - CMakeLists.txt
      - README.md
      - README_en.md
      - add_custom_kernel
        add_custom.asc
        add_custom_kernel.h
      - ge
        add_custom_ir.h
        custom_op.cpp
        utils
        log.h
        rtc_kernel_loader.cpp
        rtc_kernel_loader.h
      - run.sh
      - session_run
        main.cc
    - python
      - CMakeLists.txt
      - README.md
      - README_en.md
      - proto
        add_custom.cc
        add_custom.h
      - run.sh
      - src
        ge
        annotated_add_custom.py
        run.py
  - python
    - README.md
    - README_en.md
    - src / ge
      - annotated_add_custom.py
- args_refresh_add_custom / python
  - CMakeLists.txt
  - README.md
  - README_en.md
  - run.sh
  - src / ge
    - add_custom.py
runtime / v1 / executor
- ge_executor.cc

共 76 个文件变更+1607-923

Mdocs/en/design/modules/ge_python/ge_python_custom_op_design.md+17-8

	@@ -355,7 +355,7 @@ Python custom op loading is managed by `runtime/custom_op` to avoid direct Pytho
355	- The bridge imports `_ge_custom_op_native` and `ge.custom_op._bridge` and obtains one prototype/implementation snapshot. It registers all prototypes first, then validates and registers Adapters.	355	- The bridge imports `_ge_custom_op_native` and `ge.custom_op._bridge` and obtains one prototype/implementation snapshot. It registers all prototypes first, then validates and registers Adapters.
356	- `CustomOpLoader::LoadCustomOps()` records whether Python custom ops are loaded, so repeated lifecycle load requests return success without invoking the bridge registration entry again. The dynamic `LoadPythonCustomOpsIfNeeded()` path intentionally does not use this flag, allowing newly added Python custom op paths to be discovered during runtime. The lower-level `LoadPythonCustomOps()` function performs one bridge registration attempt; callers are responsible for invoking `UnloadPythonCustomOps()` after a failed attempt so that partial registrations are cleaned up.	356	- `CustomOpLoader::LoadCustomOps()` records whether Python custom ops are loaded, so repeated lifecycle load requests return success without invoking the bridge registration entry again. The dynamic `LoadPythonCustomOpsIfNeeded()` path intentionally does not use this flag, allowing newly added Python custom op paths to be discovered during runtime. The lower-level `LoadPythonCustomOps()` function performs one bridge registration attempt; callers are responsible for invoking `UnloadPythonCustomOps()` after a failed attempt so that partial registrations are cleaned up.
357	- `UnloadPythonCustomOps()` removes registered Adapter creators, clears the Python custom-op runtime registry in one operation, and then removes registered proto creators. It does not perform per-entry runtime unregistration or maintain pending-cleanup state in the bridge loader.	357	- `UnloadPythonCustomOps()` removes registered Adapter creators, clears the Python custom-op runtime registry in one operation, and then removes registered proto creators. It does not perform per-entry runtime unregistration or maintain pending-cleanup state in the bridge loader.
358	-- `UnloadCustomOps()` uses an `active_users_` reference count to manage the lifecycle: each `LoadCustomOps()` increments the count by 1, each `UnloadCustomOps()` decrements it by 1, and Python custom ops are only unloaded (holders/registry cleaned up and bridge closed) when the count reaches zero. `ShutdownCustomOpsForProcess()` is retained as a compatibility wrapper that internally calls `UnloadCustomOps()`.	358	+- `UnloadCustomOps()` uses an `active_users_` reference count to manage the lifecycle: each `LoadCustomOps()` increments the count by 1, each `UnloadCustomOps()` decrements it by 1, and Python custom ops are only unloaded (holders/registry cleaned up and bridge closed) when the count reaches zero.
359		359
360	Output	360	Output
361		361
	@@ -375,7 +375,7 @@ Python custom op loading is managed by `runtime/custom_op` to avoid direct Pytho
375		375
376	- Python UT covers plain-class registration, reflection of both capabilities, schema-bound signature validation during descriptor loading, schema-bound invocation, `declare_launch_args` return validation, flattened instance indices, context scope, holder lifecycle, and environment-variable plugin loading.	376	- Python UT covers plain-class registration, reflection of both capabilities, schema-bound signature validation during descriptor loading, schema-bound invocation, `declare_launch_args` return validation, flattened instance indices, context scope, holder lifecycle, and environment-variable plugin loading.
377	- C++ UT should cover the capability helper, canonical IR lookup, adapter execute/declare forwarding, loader skip and load paths, bridge ABI v1 verification, and shutdown order.	377	- C++ UT should cover the capability helper, canonical IR lookup, adapter execute/declare forwarding, loader skip and load paths, bridge ABI v1 verification, and shutdown order.
378	-- The sample `examples/custom_op/annotated_args_refresh_add_custom/python` verifies offline compilation, loading without Python, and address-refresh execution.	378	+- The samples `examples/custom_op/annotated_args_refresh_add_custom/{online,offline}/python` verify public Python `register_op`/`infer_meta`/`register_op_impl` with the Ascend C kernel online and offline. Both online operators are Python implementations; `AnnotatedAddCustom` has no C++ creator in the Python sample.
379		379
380	#### 3.3.3 Portability	380	#### 3.3.3 Portability
381		381
	@@ -416,11 +416,20 @@ The current implementation does not serialize Python implementations into the OM
416		416
417	The Python `execute` path enters the Python GIL and calls back user Python code. Its performance is not equivalent to that of C++ custom ops. The schema-bound form also traverses IR inputs and attributes and creates Python `list` / `dict` arguments, with cost growing linearly with the number of prototype parameters. This interface primarily provides development convenience and host-side scheduling and is not suitable as an ultimate execution performance path. The execution hot path does not print high-frequency logs. Logs, dynamic allocation, ACL calls, and kernel args management in user Python code are controlled by the user.	417	The Python `execute` path enters the Python GIL and calls back user Python code. Its performance is not equivalent to that of C++ custom ops. The schema-bound form also traverses IR inputs and attributes and creates Python `list` / `dict` arguments, with cost growing linearly with the number of prototype parameters. This interface primarily provides development convenience and host-side scheduling and is not suitable as an ultimate execution performance path. The execution hot path does not print high-frequency logs. Logs, dynamic allocation, ACL calls, and kernel args management in user Python code are controlled by the user.
418		418
		419	+### 4.4 Benchmark Protocol
		420	+
		421	+On an NPU build, record the median and p95 of a single-node compile and RT2
		422	+shape-inference run separately from kernel execution. Run the same graph once
		423	+with a no-op infer-meta callback and once with the real callback, using the
		424	+same Python/runtime artifact set and a fresh process for each case. Report the
		425	+callback-inclusive delta and do not mix artifacts from different Python minor
		426	+versions or builds.
		427	+
419	## 5. Interface Design	428	## 5. Interface Design
420		429
421	### 5.1 New/Modified Interface Description	430	### 5.1 New/Modified Interface Description
422		431
423	-For the Python external API, refer to `docs/zh/api/graph_engine_api/python/ge/custom_op/`. The current public interfaces are as follows:	432	+The current public Python interfaces are as follows:
424		433
425	\| Interface \| Description \|	434	\| Interface \| Description \|
426	\|-----------\|-------------\|	435	\|-----------\|-------------\|
	@@ -436,14 +445,14 @@ For the Python external API, refer to `docs/zh/api/graph_engine_api/python/ge/cu
436	\| `get_registered_op_impl_by_descriptor_key` \| Queries a descriptor by descriptor key \|	445	\| `get_registered_op_impl_by_descriptor_key` \| Queries a descriptor by descriptor key \|
437	\| `clear_registered_op_impls` \| Clears the Python registry \|	446	\| `clear_registered_op_impls` \| Clears the Python registry \|
438		447
439	-`Tensor`, `Shape`, `StorageShape`, `StorageFormat`, and `TensorPlacement` in `ge.runtime` are input and return types of the context and are not included in the `__all__` of `ge.custom_op`.	448	+`TensorDesc` is publicly available from `ge.runtime` and is the input/output type of `register_op` infer-meta functions. `Tensor`, `Shape`, `StorageShape`, `StorageFormat`, and `TensorPlacement` in `ge.runtime` are context types and are not included in the `__all__` of `ge.custom_op`.
440		449
441	### 5.2 Interface Check Items	450	### 5.2 Interface Check Items
442		451
443	\| Check Item \| Sub-Check Item \| Involved \|	452	\| Check Item \| Sub-Check Item \| Involved \|
444	\|------------\|----------------\|----------\|	453	\|------------\|----------------\|----------\|
445	\| Interface description \| Whether review is required; review should focus on interface compatibility and interface constraints \| Involved; new Python external API added \|	454	\| Interface description \| Whether review is required; review should focus on interface compatibility and interface constraints \| Involved; new Python external API added \|
446	-\| Interface description \| Whether supplementary materials are needed \| Involved; API documentation and samples have been added \|	455	+\| Interface description \| Whether supplementary materials are needed \| Involved; samples and design documentation have been added \|
447	\| Interface description \| Whether interface prototypes, functions, and return values are clearly described \| Involved; refer to the API documentation \|	456	\| Interface description \| Whether interface prototypes, functions, and return values are clearly described \| Involved; refer to the API documentation \|
448	\| Interface compatibility \| Whether behavior changes before and after modification \| Involved; existing inheritance is preserved, while execution uses schema-bound invocation \|	457	\| Interface compatibility \| Whether behavior changes before and after modification \| Involved; existing inheritance is preserved, while execution uses schema-bound invocation \|
449	\| Interface compatibility \| Whether the new interface works properly on older versions \| Involved; it is unavailable when the older run package lacks the corresponding native/bridge \|	458	\| Interface compatibility \| Whether the new interface works properly on older versions \| Involved; it is unavailable when the older run package lacks the corresponding native/bridge \|
	@@ -644,8 +653,8 @@ The implementation follows the existing Python pass and GE runtime style:
644		653
645	- Python API test entries: `ge.custom_op`, `ge.custom_op.proto`, `ge.custom_op._bridge`, `ge.custom_op.bootstrap`.	654	- Python API test entries: `ge.custom_op`, `ge.custom_op.proto`, `ge.custom_op._bridge`, `ge.custom_op.bootstrap`.
646	- Native context test entries: Eager/Compile/AnnotatedArgs borrowed contexts, `AnnotatedKernelArgs`, and launch-info methods.	655	- Native context test entries: Eager/Compile/AnnotatedArgs borrowed contexts, `AnnotatedKernelArgs`, and launch-info methods.
647	-- C++ test entries: `CustomOpCast<T>`, `PythonCustomOpAdapter`, `AnnotatedKernelArgs`, `CustomTaskInfo`, `LoadPythonCustomOps()`, `LoadCustomOps()`/`UnloadCustomOps()`, and `ShutdownCustomOpsForProcess()`.	656	+- C++ test entries: `CustomOpCast<T>`, `PythonCustomOpAdapter`, `AnnotatedKernelArgs`, `CustomTaskInfo`, `LoadPythonCustomOps()`, and `LoadCustomOps()`/`UnloadCustomOps()`.
648	-- End-to-end sample entry: `examples/custom_op/annotated_args_refresh_add_custom/python/run.sh`.	657	+- End-to-end sample entries: `examples/custom_op/annotated_args_refresh_add_custom/online/python/run.sh` and `examples/custom_op/annotated_args_refresh_add_custom/offline/python/run.sh`.
649		658
650	### 9.2 Test Design	659	### 9.2 Test Design
651		660
	@@ -664,7 +673,7 @@ The implementation follows the existing Python pass and GE runtime style:
664	\| Function \| Loader skips when no Python entry exists and loads the bridge when entries exist \| C++ gtest / stub \| UT \|	673	\| Function \| Loader skips when no Python entry exists and loads the bridge when entries exist \| C++ gtest / stub \| UT \|
665	\| Compatibility \| C++ custom op bare capability inheritance still casts correctly \| C++ gtest \| UT \|	674	\| Compatibility \| C++ custom op bare capability inheritance still casts correctly \| C++ gtest \| UT \|
666	\| Feature cross \| Ops kernel info is refreshed after online PreRun loading \| GE graph execution related tests \| UT/ST \|	675	\| Feature cross \| Ops kernel info is refreshed after online PreRun loading \| GE graph execution related tests \| UT/ST \|
667	-\| Sample \| Python custom op schema-bound graph execution plus AnnotatedArgs address refresh after offline compilation without a Python runtime environment \| `annotated_args_refresh_add_custom/python` \| ST/hardware \|	676	+\| Sample \| Python `register_op`/`infer_meta`/`register_op_impl`, Ascend C kernel schema-bound online execution, and AnnotatedArgs address refresh after offline compilation without a Python runtime environment \| `annotated_args_refresh_add_custom/{online,offline}/python` \| ST/hardware \|
668		677
669	### 9.3 Test Framework Design	678	### 9.3 Test Framework Design
670		679

Mdocs/zh/api/graph_engine_api/README.md+22-0

	@@ -939,6 +939,28 @@
939	- [Pyatc接口](python/ge/pyatc/pyatc_interface.md)	939	- [Pyatc接口](python/ge/pyatc/pyatc_interface.md)
940	- [Pyatc](python/ge/pyatc/Pyatc.md)	940	- [Pyatc](python/ge/pyatc/Pyatc.md)
941		941
		942	+ - [运行时数据结构](python/ge/runtime/overview.md)
		943	+ - [Tensor](python/ge/runtime/Tensor/Tensor.md)
		944	+ - [简介](python/ge/runtime/Tensor/overview.md)
		945	+ - [addr](python/ge/runtime/Tensor/addr.md)
		946	+ - [data\_type](python/ge/runtime/Tensor/data_type.md)
		947	+ - [expand\_dims\_type](python/ge/runtime/Tensor/expand_dims_type.md)
		948	+ - [format](python/ge/runtime/Tensor/format.md)
		949	+ - [origin\_format](python/ge/runtime/Tensor/origin_format.md)
		950	+ - [origin\_shape](python/ge/runtime/Tensor/origin_shape.md)
		951	+ - [placement](python/ge/runtime/Tensor/placement.md)
		952	+ - [shape](python/ge/runtime/Tensor/shape.md)
		953	+ - [shape\_size](python/ge/runtime/Tensor/shape_size.md)
		954	+ - [size](python/ge/runtime/Tensor/size.md)
		955	+ - [storage\_format](python/ge/runtime/Tensor/storage_format.md)
		956	+ - [storage\_shape](python/ge/runtime/Tensor/storage_shape.md)
		957	+
		958	+ - [TensorDesc](python/ge/runtime/TensorDesc/TensorDesc.md)
		959	+ - [简介](python/ge/runtime/TensorDesc/overview.md)
		960	+ - [TensorDesc构造函数](python/ge/runtime/TensorDesc/TensorDesc_constructor.md)
		961	+ - [data\_type](python/ge/runtime/TensorDesc/data_type.md)
		962	+ - [shape](python/ge/runtime/TensorDesc/shape.md)
		963	+
942	- [图基础数据结构和接口](python/ge/graph_basic_data_structure_and_interface.md)	964	- [图基础数据结构和接口](python/ge/graph_basic_data_structure_and_interface.md)
943	- [Allocator](python/ge/allocator/Allocator/Allocator.md)	965	- [Allocator](python/ge/allocator/Allocator/Allocator.md)
944	- [简介](python/ge/allocator/Allocator/overview.md)	966	- [简介](python/ge/allocator/Allocator/overview.md)

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/Tensor.md+3-0

@@ -0,0 +1,3 @@
	1	+# Tensor
	Sophia121324 天前新增的这些接口，没有修改外层的README.md L lfz2812 24 天前评论： 24 天前评论：
	2	+
	3	+`Tensor`是Python自定义算子执行回调期间由GE运行时提供的张量视图。

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/addr.md+35-0

@@ -0,0 +1,35 @@
	1	+# addr
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取Tensor数据在运行时的地址。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+tensor.addr -> int
	15	+```
	16	+
	17	+## 参数说明
	18	+
	19	+无
	20	+
	21	+## 返回值说明
	22	+
	23	+返回Tensor数据地址的整数表示。
	24	+
	25	+## 约束说明
	26	+
	27	+- `Tensor`由运行时提供，只能在当前执行回调期间使用。
	28	+- 该地址仅用于运行时执行参数构造，不应在Python中对其进行内存读写或释放。
	29	+
	30	+## 调用示例
	31	+
	32	+```python
	33	+def execute(self, x: Tensor, y: Tensor) -> None:
	34	+ address = x.addr
	35	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/data_type.md+33-0

@@ -0,0 +1,33 @@
	1	+# data\_type
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取Tensor的数据类型。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+tensor.data_type -> DataType
	15	+```
	16	+
	17	+## 参数说明
	18	+
	19	+无
	20	+
	21	+## 返回值说明
	22	+
	23	+返回Tensor的`DataType`。
	24	+
	25	+## 约束说明
	26	+
	27	+返回值为DataType枚举值的拷贝。
	28	+
	29	+## 调用示例
	30	+
	31	+```python
	32	+data_type = x.data_type
	33	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/expand_dims_type.md+34-0

@@ -0,0 +1,34 @@
	1	+# expand\_dims\_type
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取OriginShape转换为StorageShape时使用的维度扩展规则。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+tensor.expand_dims_type -> ExpandDimsType
	15	+```
	16	+
	17	+## 参数说明
	18	+
	19	+无
	20	+
	21	+## 返回值说明
	22	+
	23	+返回当前Tensor Format中的`ExpandDimsType`对象。该对象由运行时提供，仅在当前执行回调期间有效。
	24	+
	25	+## 约束说明
	26	+
	27	+- 返回对象只能在当前执行回调期间使用。
	28	+- 该规则描述格式引入的扩展维度，不表示数据拷贝或内存分配操作。
	29	+
	30	+## 调用示例
	31	+
	32	+```python
	33	+expand_dims_type = x.expand_dims_type
	34	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/format.md+34-0

@@ -0,0 +1,34 @@
	1	+# format
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取Tensor的Format描述，其中同时保存OriginFormat、StorageFormat和维度扩展规则。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+tensor.format -> StorageFormat
	15	+```
	16	+
	17	+## 参数说明
	18	+
	19	+无
	20	+
	21	+## 返回值说明
	22	+
	23	+返回当前Tensor的`StorageFormat`对象。该对象由运行时提供，仅在当前执行回调期间有效。
	24	+
	25	+## 约束说明
	26	+
	27	+- 返回对象只能在当前执行回调期间使用。
	28	+- 申请输出Tensor时，可将该对象传入`EagerOpExecutionContext.malloc_output_tensor`。
	29	+
	30	+## 调用示例
	31	+
	32	+```python
	33	+storage_format = x.format
	34	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/origin_format.md+33-0

@@ -0,0 +1,33 @@
	1	+# origin\_format
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取Tensor的原始Format。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+tensor.origin_format -> Format
	15	+```
	16	+
	17	+## 参数说明
	18	+
	19	+无
	20	+
	21	+## 返回值说明
	22	+
	23	+返回表示原始逻辑布局的`Format`。
	24	+
	25	+## 约束说明
	26	+
	27	+返回值是Format枚举值的拷贝，可以在当前回调外保存；但不应据此延长Tensor的使用生命周期。
	28	+
	29	+## 调用示例
	30	+
	31	+```python
	32	+origin_format = x.origin_format
	33	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/origin_shape.md+34-0

@@ -0,0 +1,34 @@
	1	+# origin\_shape
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取Tensor的原始逻辑Shape。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+tensor.origin_shape -> Shape
	15	+```
	16	+
	17	+## 参数说明
	18	+
	19	+无
	20	+
	21	+## 返回值说明
	22	+
	23	+返回表示原始逻辑形状的`Shape`对象。该对象由运行时提供，仅在当前执行回调期间有效。
	24	+
	25	+## 约束说明
	26	+
	27	+- 返回对象只能在当前执行回调期间使用。
	28	+- 需要运行时实际内存布局时，应使用[`storage_shape`](storage_shape.md)。
	29	+
	30	+## 调用示例
	31	+
	32	+```python
	33	+origin_shape = x.origin_shape
	34	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/overview.md+22-0

@@ -0,0 +1,22 @@
	1	+# 简介
	2	+
	3	+`ge.runtime.Tensor`是Python自定义算子执行回调期间使用的张量运行时视图。它提供张量地址、大小、Shape、Format、DataType和存储位置等元数据，不负责分配、释放或拷贝张量数据。
	4	+
	5	+`Tensor`由GE运行时通过`execute`、`declare_launch_args`或执行context提供，用户不能直接构造。由回调返回的`Tensor`及其`Shape`、`StorageShape`、`StorageFormat`和`ExpandDimsType`视图，仅在当前回调有效。
	6	+
	7	+下面的示例在`execute`回调中读取输入Tensor的元数据，并按照输入Tensor的Shape、Format和DataType申请输出Tensor。
	8	+
	9	+```python
	10	+from ge.custom_op import get_execute_ctx, register_op_impl
	11	+from ge.runtime import Tensor
	12	+
	13	+
	14	+@register_op_impl(op_type="AddPythonCustomOp")
	15	+class AddPythonCustomOp:
	16	+ def execute(self, x: Tensor, y: Tensor) -> None:
	17	+ ctx = get_execute_ctx()
	18	+ output = ctx.malloc_output_tensor(0, x.shape, x.format, x.data_type)
	19	+ print(x.addr, x.size, x.storage_shape.dims, x.placement)
	20	+ # output 是当前执行回调中的 Tensor，可继续用于执行参数构造。
	21	+ _ = output
	22	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/placement.md+33-0

@@ -0,0 +1,33 @@
	1	+# placement
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取Tensor数据的存储位置。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+tensor.placement -> TensorPlacement
	15	+```
	16	+
	17	+## 参数说明
	18	+
	19	+无
	20	+
	21	+## 返回值说明
	22	+
	23	+返回`TensorPlacement`，用于表示数据位于Device HBM、Host、Following或Device P2P等位置。
	24	+
	25	+## 约束说明
	26	+
	27	+返回值为TensorPlacement枚举值的拷贝。
	28	+
	29	+## 调用示例
	30	+
	31	+```python
	32	+placement = x.placement
	33	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/shape.md+34-0

@@ -0,0 +1,34 @@
	1	+# shape
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取Tensor的Shape描述，其中同时保存OriginShape和StorageShape。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+tensor.shape -> StorageShape
	15	+```
	16	+
	17	+## 参数说明
	18	+
	19	+无
	20	+
	21	+## 返回值说明
	22	+
	23	+返回当前Tensor的`StorageShape`对象。该对象由运行时提供，仅在当前执行回调期间有效。
	24	+
	25	+## 约束说明
	26	+
	27	+- 返回对象只能在当前执行回调期间使用。
	28	+- 返回对象的`origin_shape`表示原始逻辑形状，`storage_shape`表示运行时存储形状。
	29	+
	30	+## 调用示例
	31	+
	32	+```python
	33	+storage_shape = x.shape
	34	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/shape_size.md+33-0

@@ -0,0 +1,33 @@
	1	+# shape\_size
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取Tensor的StorageShape所表示的元素个数。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+tensor.shape_size -> int
	15	+```
	16	+
	17	+## 参数说明
	18	+
	19	+无
	20	+
	21	+## 返回值说明
	22	+
	23	+返回StorageShape的元素个数。
	24	+
	25	+## 约束说明
	26	+
	27	+该值根据StorageShape计算，不能用来替代`size`取得字节数。
	28	+
	29	+## 调用示例
	30	+
	31	+```python
	32	+element_count = x.shape_size
	33	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/size.md+33-0

@@ -0,0 +1,33 @@
	1	+# size
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取Tensor数据占用的内存大小。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+tensor.size -> int
	15	+```
	16	+
	17	+## 参数说明
	18	+
	19	+无
	20	+
	21	+## 返回值说明
	22	+
	23	+返回Tensor数据占用的字节数。
	24	+
	25	+## 约束说明
	26	+
	27	+返回值表示运行时存储空间大小，不一定等于OriginShape的元素个数乘以单个元素字节数。
	28	+
	29	+## 调用示例
	30	+
	31	+```python
	32	+byte_size = x.size
	33	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/storage_format.md+33-0

@@ -0,0 +1,33 @@
	1	+# storage\_format
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取Tensor的运行时存储Format。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+tensor.storage_format -> Format
	15	+```
	16	+
	17	+## 参数说明
	18	+
	19	+无
	20	+
	21	+## 返回值说明
	22	+
	23	+返回表示实际内存布局的`Format`。
	24	+
	25	+## 约束说明
	26	+
	27	+返回值是Format枚举值的拷贝，可以在当前回调外保存；但不应据此延长Tensor的使用生命周期。
	28	+
	29	+## 调用示例
	30	+
	31	+```python
	32	+storage_format = x.storage_format
	33	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/Tensor/storage_shape.md+34-0

@@ -0,0 +1,34 @@
	1	+# storage\_shape
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取Tensor的运行时存储Shape。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+tensor.storage_shape -> Shape
	15	+```
	16	+
	17	+## 参数说明
	18	+
	19	+无
	20	+
	21	+## 返回值说明
	22	+
	23	+返回表示实际存储形状的`Shape`对象。该对象由运行时提供，仅在当前执行回调期间有效。
	24	+
	25	+## 约束说明
	26	+
	27	+- 返回对象只能在当前执行回调期间使用。
	28	+- `storage_shape`可能因格式转换或对齐而与[`origin_shape`](origin_shape.md)不同。
	29	+
	30	+## 调用示例
	31	+
	32	+```python
	33	+storage_shape = x.storage_shape
	34	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/TensorDesc/TensorDesc.md+9-0

@@ -0,0 +1,9 @@
	1	+# TensorDesc
	2	+
	3	+`TensorDesc`是Python自定义算子`infer_meta`使用的张量元信息描述。
	4	+
	5	+## 接口列表
	6	+
	7	+- [`TensorDesc构造函数`](TensorDesc_constructor.md)：创建TensorDesc。
	8	+- [`data_type`](data_type.md)：获取或设置数据类型。
	9	+- [`shape`](shape.md)：获取或设置StorageShape。

Adocs/zh/api/graph_engine_api/python/ge/runtime/TensorDesc/TensorDesc_constructor.md+40-0

@@ -0,0 +1,40 @@
	1	+# TensorDesc构造函数
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+创建一个用于Python自定义算子元信息推导的`TensorDesc`。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+TensorDesc(shape: Optional[Union[StorageShape, List[int]]], data_type: DataType)
	15	+```
	16	+
	17	+## 参数说明
	18	+
	19	+\| 参数名 \| 输入/输出 \| 描述 \|
	20	+\| --- \| --- \| --- \|
	21	+\| shape \| 输入 \| Tensor的Shape，类型为`StorageShape`或`List[int]`。`None`表示标量。 \|
	22	+\| data_type \| 输入 \| Tensor的数据类型，类型为`DataType`。 \|
	23	+
	24	+## 返回值说明
	25	+
	26	+TensorDesc对象。
	27	+
	28	+## 约束说明
	29	+
	30	+- `shape`不是`StorageShape`、整数列表或`None`时抛出`TypeError`。
	31	+- `data_type`不是`DataType`或等于`DataType.DT_MAX`时抛出`TypeError`或`ValueError`。
	32	+
	33	+## 调用示例
	34	+
	35	+```python
	36	+from ge.graph import DataType
	37	+from ge.runtime import TensorDesc
	38	+
	39	+desc = TensorDesc([2, 3], DataType.DT_FLOAT)
	40	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/TensorDesc/data_type.md+38-0

@@ -0,0 +1,38 @@
	1	+# data\_type
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取或设置TensorDesc的DataType。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+desc.data_type -> DataType
	15	+desc.data_type = data_type
	16	+```
	17	+
	18	+## 参数说明
	19	+
	20	+\| 参数名 \| 输入/输出 \| 描述 \|
	21	+\| --- \| --- \| --- \|
	22	+\| data\_type \| 输入 \| Tensor的数据类型，类型为DataType。 \|
	23	+
	24	+## 返回值说明
	25	+
	26	+获取时返回DataType；设置时无返回值。
	27	+
	28	+## 约束说明
	29	+
	30	+- `data_type`不是DataType类型时，抛出TypeError。
	31	+- `data_type`等于DataType.DT_MAX时，抛出ValueError。
	32	+
	33	+## 调用示例
	34	+
	35	+```python
	36	+data_type = desc.data_type
	37	+desc.data_type = DataType.DT_FLOAT16
	38	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/TensorDesc/overview.md+23-0

@@ -0,0 +1,23 @@
	1	+# 简介
	2	+
	3	+`TensorDesc`类用于表示Python自定义算子`infer_meta`的张量元信息，保存Tensor的逻辑Shape、StorageShape和DataType。该对象由Python代码持有，可以作为`infer_meta`的返回值。
	4	+
	5	+注意：该类型与[`ge.graph.TensorDesc`](../../graph/TensorDesc/overview.md)不同。`ge.graph.TensorDesc`用于图构建侧的Tensor描述，包含Format、OriginShape以及对应的`get_`/`set_`接口；`ge.runtime.TensorDesc`仅用于Python原型的元信息输入和输出。
	6	+
	7	+以下示例创建两个输入TensorDesc，并调用`infer_meta`返回输出TensorDesc。
	8	+
	9	+```python
	10	+from ge.custom_op import register_op
	11	+from ge.graph import DataType
	12	+from ge.runtime import TensorDesc
	13	+
	14	+
	15	+@register_op(op_type="AddCustom")
	16	+def infer_meta(x: TensorDesc, y: TensorDesc) -> TensorDesc:
	17	+ return TensorDesc(x.shape, x.data_type)
	18	+
	19	+
	20	+x = TensorDesc([2, 3], DataType.DT_FLOAT)
	21	+y = TensorDesc([2, 3], DataType.DT_FLOAT)
	22	+z = infer_meta(x, y)
	23	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/TensorDesc/shape.md+38-0

@@ -0,0 +1,38 @@
	1	+# shape
	2	+
	3	+## 产品支持情况
	4	+
	5	+全量芯片支持。
	6	+
	7	+## 功能说明
	8	+
	9	+获取或设置TensorDesc的StorageShape。
	10	+
	11	+## 函数原型
	12	+
	13	+```python
	14	+desc.shape -> StorageShape
	15	+desc.shape = shape
	16	+```
	17	+
	18	+## 参数说明
	19	+
	20	+\| 参数名 \| 输入/输出 \| 描述 \|
	21	+\| --- \| --- \| --- \|
	22	+\| shape \| 输入 \| Tensor的Shape，类型为`StorageShape`或`List[int]`。 \|
	23	+
	24	+## 返回值说明
	25	+
	26	+获取时返回StorageShape；设置时无返回值。
	27	+
	28	+## 约束说明
	29	+
	30	+- `shape`不是StorageShape、整数列表或None时，抛出TypeError。
	31	+- 返回的StorageShape隶属于当前TensorDesc对象，TensorDesc对象销毁后不可继续使用。
	32	+
	33	+## 调用示例
	34	+
	35	+```python
	36	+shape = desc.shape
	37	+desc.shape = [2, 3]
	38	+```

Adocs/zh/api/graph_engine_api/python/ge/runtime/overview.md+8-0

@@ -0,0 +1,8 @@
	1	+# 简介
	2	+
	3	+`ge.runtime`提供Python自定义算子执行和元信息推导使用的运行时数据结构。`TensorDesc`用于`register_op`装饰函数的输入和返回值；`Tensor`、`StorageShape`、`StorageFormat`等对象由运行时context在回调期间提供。
	4	+
	5	+其中`ge.runtime.TensorDesc`与图构建API中的[`ge.graph.TensorDesc`](../graph/TensorDesc/overview.md)是不同类型，使用场景和接口不可混用。
	6	+
	7	+- [`Tensor`](Tensor/overview.md)：执行回调期间使用的张量运行时视图。
	8	+- [`TensorDesc`](TensorDesc/overview.md)：Python自定义算子`infer_meta`使用的张量元信息。

Mdocs/zh/design/modules/ge_python/ge_python_custom_op_design.md+13-6

	@@ -355,7 +355,7 @@ Python custom op 加载由 `runtime/custom_op` 管理，避免 `graph_metadef/re
355	- bridge 导入 `_ge_custom_op_native` 和 `ge.custom_op._bridge`，一次获取 proto/impl snapshot；先注册全部 proto，再校验并注册 Adapter。	355	- bridge 导入 `_ge_custom_op_native` 和 `ge.custom_op._bridge`，一次获取 proto/impl snapshot；先注册全部 proto，再校验并注册 Adapter。
356	- `CustomOpLoader::LoadCustomOps()` 记录 Python custom op 是否已经加载，因此生命周期加载请求重复调用时会直接返回成功，不重复调用 bridge 注册入口。动态 `LoadPythonCustomOpsIfNeeded()` 路径不使用该状态，运行期间可以继续发现新增加的 Python custom op 路径。底层 `LoadPythonCustomOps()` 负责执行一次 bridge 注册尝试；注册失败后由调用方调用 `UnloadPythonCustomOps()` 清理本次产生的部分注册。	356	- `CustomOpLoader::LoadCustomOps()` 记录 Python custom op 是否已经加载，因此生命周期加载请求重复调用时会直接返回成功，不重复调用 bridge 注册入口。动态 `LoadPythonCustomOpsIfNeeded()` 路径不使用该状态，运行期间可以继续发现新增加的 Python custom op 路径。底层 `LoadPythonCustomOps()` 负责执行一次 bridge 注册尝试；注册失败后由调用方调用 `UnloadPythonCustomOps()` 清理本次产生的部分注册。
357	- `UnloadPythonCustomOps()` 先移除已注册的 Adapter creator，再一次性清理 Python 自定义算子 runtime registry，最后清理已注册的 proto creator。bridge loader 不再逐项注销 runtime entry，也不维护待清理状态。	357	- `UnloadPythonCustomOps()` 先移除已注册的 Adapter creator，再一次性清理 Python 自定义算子 runtime registry，最后清理已注册的 proto creator。bridge loader 不再逐项注销 runtime entry，也不维护待清理状态。
358	-- `UnloadCustomOps()` 采用 `active_users_` 引用计数管理生命周期：每次 `LoadCustomOps()` 使计数 +1，每次 `UnloadCustomOps()` 使计数 -1，仅当计数归零时才卸载 Python custom op、清理 Python holder/registry 并关闭 bridge。`ShutdownCustomOpsForProcess()` 作为兼容 wrapper 保留，内部调用 `UnloadCustomOps()`。	358	+- `UnloadCustomOps()` 采用 `active_users_` 引用计数管理生命周期：每次 `LoadCustomOps()` 使计数 +1，每次 `UnloadCustomOps()` 使计数 -1，仅当计数归零时才卸载 Python custom op、清理 Python holder/registry 并关闭 bridge。
359		359
360	输出	360	输出
361		361
	@@ -375,7 +375,7 @@ Python custom op 加载由 `runtime/custom_op` 管理，避免 `graph_metadef/re
375		375
376	- Python UT 覆盖普通 class 注册、两类能力反射、descriptor 加载阶段的 schema-bound 签名校验、schema-bound 调用、`declare_launch_args` 返回值校验、实例平铺 index、context 作用域、holder 生命周期和环境变量插件加载。	376	- Python UT 覆盖普通 class 注册、两类能力反射、descriptor 加载阶段的 schema-bound 签名校验、schema-bound 调用、`declare_launch_args` 返回值校验、实例平铺 index、context 作用域、holder 生命周期和环境变量插件加载。
377	- C++ UT 应覆盖 capability helper、canonical IR 查询、adapter 的 execute/declare 转发、loader 跳过/加载路径、bridge ABI v1 校验和 shutdown 顺序。	377	- C++ UT 应覆盖 capability helper、canonical IR 查询、adapter 的 execute/declare 转发、loader 跳过/加载路径、bridge ABI v1 校验和 shutdown 顺序。
378	-- 样例 `examples/custom_op/annotated_args_refresh_add_custom/python` 验证离线编译、无 Python 环境加载和地址刷新执行。	378	+- 样例 `examples/custom_op/annotated_args_refresh_add_custom/{online,offline}/python` 分别验证公开 Python `register_op`/`infer_meta`/`register_op_impl` 与 Ascend C kernel 的在线、离线链路；两个在线算子均由 Python 实现，AnnotatedAddCustom 不再与 C++ creator 共存。
379		379
380	#### 3.3.3 可移植性	380	#### 3.3.3 可移植性
381		381
	@@ -416,6 +416,13 @@ Python custom op 不区分芯片，不引入芯片分支。device kernel 能力
416		416
417	Python `execute` 路径会进入 Python GIL，并回调用户 Python 代码，性能不等同于 C++ custom op。schema-bound 形式还会按 IR 遍历输入和属性并创建 Python `list` / `dict` 实参，成本随原型参数数量线性增长。该接口主要用于开发便利性和 host 侧调度能力，不适合作为极致执行性能路径。执行热路径不额外打印高频日志；用户 Python 代码中的日志、动态分配、ACL 调用和 kernel args 管理由用户自行控制。	417	Python `execute` 路径会进入 Python GIL，并回调用户 Python 代码，性能不等同于 C++ custom op。schema-bound 形式还会按 IR 遍历输入和属性并创建 Python `list` / `dict` 实参，成本随原型参数数量线性增长。该接口主要用于开发便利性和 host 侧调度能力，不适合作为极致执行性能路径。执行热路径不额外打印高频日志；用户 Python 代码中的日志、动态分配、ACL 调用和 kernel args 管理由用户自行控制。
418		418
		419	+### 4.4 基准记录方法
		420	+
		421	+在 NPU 构建环境分别记录单节点编译和 RT2 shape 推导的 median、p95，
		422	+并将其与 kernel 执行耗时分开。相同图分别使用空 infer-meta 回调和真实回调，
		423	+每种情况使用同一组 Python/runtime artifact 并在新进程中执行。回归值记录为
		424	+包含回调的耗时差；禁止混用不同 Python minor 版本或不同构建产物。
		425	+
419	## 5. 接口设计	426	## 5. 接口设计
420		427
421	### 5.1 新增/修改接口描述	428	### 5.1 新增/修改接口描述
	@@ -436,7 +443,7 @@ Python 对外 API 见 `docs/zh/api/graph_engine_api/python/ge/custom_op/`。当
436	\| `get_registered_op_impl_by_descriptor_key` \| 按 descriptor key 查询 descriptor \|	443	\| `get_registered_op_impl_by_descriptor_key` \| 按 descriptor key 查询 descriptor \|
437	\| `clear_registered_op_impls` \| 清理 Python registry \|	444	\| `clear_registered_op_impls` \| 清理 Python registry \|
438		445
439	-`ge.runtime` 中的 `Tensor`、`Shape`、`StorageShape`、`StorageFormat`、`TensorPlacement` 是 context 的入参/返回类型，不归入 `ge.custom_op` 的 `__all__`。	446	+`ge.runtime` 公开提供 `TensorDesc`，作为 `register_op` infer-meta 函数的输入/输出类型。`ge.runtime` 中的 `Tensor`、`Shape`、`StorageShape`、`StorageFormat`、`TensorPlacement` 是 context 的入参/返回类型，不归入 `ge.custom_op` 的 `__all__`。
440		447
441	### 5.2 接口检查项	448	### 5.2 接口检查项
442		449
	@@ -645,8 +652,8 @@ PythonCustomOpAdapter::DeclareLaunchArgs(ctx)
645		652
646	- Python API 测试入口：`ge.custom_op`、`ge.custom_op.proto`、`ge.custom_op._bridge`、`ge.custom_op.bootstrap`。	653	- Python API 测试入口：`ge.custom_op`、`ge.custom_op.proto`、`ge.custom_op._bridge`、`ge.custom_op.bootstrap`。
647	- Native context 测试入口：Eager/Compile/AnnotatedArgs borrowed context、`AnnotatedKernelArgs` 和 launch info 方法。	654	- Native context 测试入口：Eager/Compile/AnnotatedArgs borrowed context、`AnnotatedKernelArgs` 和 launch info 方法。
648	-- C++ 测试入口：`CustomOpCast<T>`、`PythonCustomOpAdapter`、`AnnotatedKernelArgs`、`CustomTaskInfo`、`LoadPythonCustomOps()`、`LoadCustomOps()`/`UnloadCustomOps()` 和 `ShutdownCustomOpsForProcess()`。	655	+- C++ 测试入口：`CustomOpCast<T>`、`PythonCustomOpAdapter`、`AnnotatedKernelArgs`、`CustomTaskInfo`、`LoadPythonCustomOps()` 和 `LoadCustomOps()`/`UnloadCustomOps()`。
649	-- 端到端样例入口：`examples/custom_op/annotated_args_refresh_add_custom/python/run.sh`。	656	+- 端到端样例入口：`examples/custom_op/annotated_args_refresh_add_custom/online/python/run.sh` 和 `examples/custom_op/annotated_args_refresh_add_custom/offline/python/run.sh`。
650		657
651	### 9.2 测试设计	658	### 9.2 测试设计
652		659
	@@ -667,7 +674,7 @@ PythonCustomOpAdapter::DeclareLaunchArgs(ctx)
667	\| 功能 \| loader 在无 Python 入口时跳过，有入口时加载 bridge \| C++ gtest / stub \| UT \|	674	\| 功能 \| loader 在无 Python 入口时跳过，有入口时加载 bridge \| C++ gtest / stub \| UT \|
668	\| 兼容性 \| C++ custom op 裸能力继承仍可正常 cast \| C++ gtest \| UT \|	675	\| 兼容性 \| C++ custom op 裸能力继承仍可正常 cast \| C++ gtest \| UT \|
669	\| 特性交叉 \| 在线 PreRun 加载后刷新 ops kernel info \| GE 图执行相关测试 \| UT/ST \|	676	\| 特性交叉 \| 在线 PreRun 加载后刷新 ops kernel info \| GE 图执行相关测试 \| UT/ST \|
670	-\| 样例 \| Python custom op schema-bound 构图执行，以及 AnnotatedArgs 离线编译后无 Python 环境的地址刷新执行 \| `annotated_args_refresh_add_custom/python` \| ST/真机 \|	677	+\| 样例 \| Python `register_op`/`infer_meta`/`register_op_impl`、Ascend C kernel 的 schema-bound 在线执行，以及 AnnotatedArgs 离线编译后无 Python 环境的地址刷新执行 \| `annotated_args_refresh_add_custom/{online,offline}/python` \| ST/真机 \|
671		678
672	### 9.3 测试框架设计	679	### 9.3 测试框架设计
673		680

Mexamples/custom_op/README.md+0-1

	@@ -13,7 +13,6 @@
13	\| `data_dependent_shape_custom` \| 数据依赖 shape 算子 \| GE \| Ascend C \| CMake编译 \| 不涉及 \| [README](data_dependent_shape_custom/README.md) \|	13	\| `data_dependent_shape_custom` \| 数据依赖 shape 算子 \| GE \| Ascend C \| CMake编译 \| 不涉及 \| [README](data_dependent_shape_custom/README.md) \|
14	\| `args_refresh_add_custom` \| ArgsUpdater 地址刷新 + MallocReadOnlyDevArgs + 性能对比 \| GE 在线执行 \| Ascend C \| RTC 运行时编译 \| 在线地址刷新性能对比 \| [README](./args_refresh_add_custom/cpp/README.md) \|	14	\| `args_refresh_add_custom` \| ArgsUpdater 地址刷新 + MallocReadOnlyDevArgs + 性能对比 \| GE 在线执行 \| Ascend C \| RTC 运行时编译 \| 在线地址刷新性能对比 \| [README](./args_refresh_add_custom/cpp/README.md) \|
15	\| `annotated_args_refresh_add_custom` \| AnnotatedArgsOp 声明式地址刷新在线场景性能对比+离线场景 \| GE 在线执行 + ATC 离线编译 \| Ascend C \| RTC 运行时编译 \| 支持在线性能对比和 OM 模型下沉 \| [README](./annotated_args_refresh_add_custom/README.md) \|	15	\| `annotated_args_refresh_add_custom` \| AnnotatedArgsOp 声明式地址刷新在线场景性能对比+离线场景 \| GE 在线执行 + ATC 离线编译 \| Ascend C \| RTC 运行时编译 \| 支持在线性能对比和 OM 模型下沉 \| [README](./annotated_args_refresh_add_custom/README.md) \|
16	-\| `args_refresh_add_custom（Python 版本）` \| Python EagerExecuteOp 执行 \| GE 在线执行 \| Ascend C \| Bisheng 预编译 \| 不涉及 \| [README](./args_refresh_add_custom/python/README.md) \|
17	\| `tilelang_add_custom` \| TileLang 算子通过 GE 入图 \| GE 原生 (Session API) \| TileLang \| TileLang 预编译产出 `.so` \| 不涉及 \| [README](./tilelang_add_custom/README.md) \|	16	\| `tilelang_add_custom` \| TileLang 算子通过 GE 入图 \| GE 原生 (Session API) \| TileLang \| TileLang 预编译产出 `.so` \| 不涉及 \| [README](./tilelang_add_custom/README.md) \|
18	\| `tilelang_add_custom_online` \| TileLang 算子在线编译 + 在线执行 \| GE 原生 (Session API) \| TileLang \| GE 编译阶段 `CompilableOp::Compile` subprocess 调用 Python 编译器 \| 不涉及 \| [README](./tilelang_add_custom_online/README.md) \|	17	\| `tilelang_add_custom_online` \| TileLang 算子在线编译 + 在线执行 \| GE 原生 (Session API) \| TileLang \| GE 编译阶段 `CompilableOp::Compile` subprocess 调用 Python 编译器 \| 不涉及 \| [README](./tilelang_add_custom_online/README.md) \|
19	\| `tilelang_add_custom_offline` \| TileLang 算子离线 OM 模型下沉 \| GE 原生 (`aclgrphBuildModel`) \| TileLang \| `CompilableOp::Compile` + `PortableOp::Serialize` 序列化到 OM \| 支持 OM 模型下沉 \| [README](./tilelang_add_custom_offline/README.md) \|	18	\| `tilelang_add_custom_offline` \| TileLang 算子离线 OM 模型下沉 \| GE 原生 (`aclgrphBuildModel`) \| TileLang \| `CompilableOp::Compile` + `PortableOp::Serialize` 序列化到 OM \| 支持 OM 模型下沉 \| [README](./tilelang_add_custom_offline/README.md) \|

Mexamples/custom_op/README_en.md+0-1

	@@ -13,7 +13,6 @@ This directory provides samples related to custom operator graph integration, co
13	\| `data_dependent_shape_custom` \| Data dependent shape operator \| GE \| Ascend C \| CMake compilation \| Not involved \| [README](data_dependent_shape_custom/README_en.md) \|	13	\| `data_dependent_shape_custom` \| Data dependent shape operator \| GE \| Ascend C \| CMake compilation \| Not involved \| [README](data_dependent_shape_custom/README_en.md) \|
14	\| `args_refresh_add_custom` \| ArgsUpdater address refresh + MallocReadOnlyDevArgs + performance comparison \| GE online execution \| Ascend C \| RTC runtime compilation \| Online address refresh performance comparison \| [README](./args_refresh_add_custom/cpp/README_en.md) \|	14	\| `args_refresh_add_custom` \| ArgsUpdater address refresh + MallocReadOnlyDevArgs + performance comparison \| GE online execution \| Ascend C \| RTC runtime compilation \| Online address refresh performance comparison \| [README](./args_refresh_add_custom/cpp/README_en.md) \|
15	\| `annotated_args_refresh_add_custom` \| AnnotatedArgsOp declarative address refresh for online performance comparison and offline model \| GE online execution + ATC offline compilation \| Ascend C \| RTC runtime compilation \| Supports online performance comparison and OM model sink \| [README](./annotated_args_refresh_add_custom/README_en.md) \|	15	\| `annotated_args_refresh_add_custom` \| AnnotatedArgsOp declarative address refresh for online performance comparison and offline model \| GE online execution + ATC offline compilation \| Ascend C \| RTC runtime compilation \| Supports online performance comparison and OM model sink \| [README](./annotated_args_refresh_add_custom/README_en.md) \|
16	-\| `args_refresh_add_custom (Python version)` \| Python EagerExecuteOp execution \| GE online execution \| Ascend C \| Bisheng pre-compilation \| Not involved \| [README](./args_refresh_add_custom/python/README.md) \|
17	\| `tilelang_add_custom` \| TileLang operator enters graph through GE \| GE native (Session API) \| TileLang \| TileLang pre-compiled `.so` \| Not involved \| [README](./tilelang_add_custom/README_en.md) \|	16	\| `tilelang_add_custom` \| TileLang operator enters graph through GE \| GE native (Session API) \| TileLang \| TileLang pre-compiled `.so` \| Not involved \| [README](./tilelang_add_custom/README_en.md) \|
18	\| `tilelang_add_custom_online` \| TileLang operator online compilation + online execution \| GE native (Session API) \| TileLang \| `CompilableOp::Compile` subprocess invokes Python compiler during GE compile phase \| Not involved \| [README](./tilelang_add_custom_online/README_en.md) \|	17	\| `tilelang_add_custom_online` \| TileLang operator online compilation + online execution \| GE native (Session API) \| TileLang \| `CompilableOp::Compile` subprocess invokes Python compiler during GE compile phase \| Not involved \| [README](./tilelang_add_custom_online/README_en.md) \|
19	\| `tilelang_add_custom_offline` \| TileLang operator offline OM model sinking \| GE native (`aclgrphBuildModel`) \| TileLang \| `CompilableOp::Compile` + `PortableOp::Serialize` to OM \| Supports OM model sinking \| [README](./tilelang_add_custom_offline/README_en.md) \|	18	\| `tilelang_add_custom_offline` \| TileLang operator offline OM model sinking \| GE native (`aclgrphBuildModel`) \| TileLang \| `CompilableOp::Compile` + `PortableOp::Serialize` to OM \| Supports OM model sinking \| [README](./tilelang_add_custom_offline/README_en.md) \|

Mexamples/custom_op/annotated_args_refresh_add_custom/README.md+6-5

	@@ -1,7 +1,8 @@
1	-# Ascend C 自定义算子声明式地址刷新样例	1	+# AnnotatedArgs 声明式地址刷新自定义算子样例
2		2
3	-本目录提供基于 `AnnotatedArgsOp::DeclareLaunchArgs` 的自定义算子声明式地址刷新在线与离线样例：	3	+本目录按在线、离线和语言组织样例，C++ 与 Python 分别提供独立实现：
4		4
5	-- [online](./online/README.md)：对比声明式地址刷新算子与非地址刷新算子的在线执行性能。	5	+- [online/cpp](./online/cpp/README.md)：C++ 在线性能对比样例。
6	-- [offline](./offline/README.md)：演示 AIR/OM 生成、`PortableOp` 序列化与反序列化、`DeclareLaunchArgs` 参数声明，以及通过 ACL 加载并执行 OM。	6	+- [online/python](./online/python/README.md)：Python 在线性能对比样例。
7	-- [python](./python/README.md)：Python 构图 + ATC 编译 + ACL 两轮 NPU 地址刷新验证。	7	+- [offline/cpp](./offline/cpp/README.md)：C++ AIR/OM 生成、AnnotatedArgs 声明和 ACL 执行样例。
		8	+- [offline/python](./offline/python/README.md)：Python 装饰器原型注册、infer-meta、地址声明以及离线 AIR/OM 执行样例。

Mexamples/custom_op/annotated_args_refresh_add_custom/README_en.md+6-5

	@@ -1,7 +1,8 @@
1	-# Ascend C Declarative Address Refresh Custom Operator Samples	1	+# AnnotatedArgs Declarative Address Refresh Samples
2		2
3	-This directory provides online and offline declarative address refresh samples based on `AnnotatedArgsOp::DeclareLaunchArgs`:	3	+Samples are organized by execution mode and language, with independent C++ and Python implementations:
4		4
5	-- [online](./online/README_en.md): Compares online execution performance between declarative address-refresh and no-refresh operators.	5	+- [online/cpp](./online/cpp/README_en.md): C++ online performance comparison.
6	-- [offline](./offline/README_en.md): Demonstrates AIR/OM generation, `PortableOp` serialization and deserialization, argument declaration through `DeclareLaunchArgs`, and OM loading and execution through ACL.	6	+- [online/python](./online/python/README_en.md): Python online performance comparison.
7	-- [python](./python/README_en.md): Python graph construction + ATC compilation + ACL two-round NPU address-refresh validation.	7	+- [offline/cpp](./offline/cpp/README_en.md): C++ AIR/OM generation, AnnotatedArgs declaration, and ACL execution.
		8	+- [offline/python](./offline/python/README_en.md): Python decorator prototype registration, infer-meta, address declaration, and offline AIR/OM execution.

Rexamples/custom_op/annotated_args_refresh_add_custom/offline/CMakeLists.txt→examples/custom_op/annotated_args_refresh_add_custom/offline/cpp/CMakeLists.txt+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/offline/README.md→examples/custom_op/annotated_args_refresh_add_custom/offline/cpp/README.md+2-2

	@@ -33,7 +33,7 @@ Compile 回调 RTC 编译 Ascend C kernel
33		33
34	## 快速运行	34	## 快速运行
35		35
36	-在 `examples/custom_op/annotated_args_refresh_add_custom/offline` 目录执行：	36	+在 `examples/custom_op/annotated_args_refresh_add_custom/offline/cpp` 目录执行：
37		37
38	```bash	38	```bash
39	bash run.sh	39	bash run.sh
	@@ -57,7 +57,7 @@ First element of output: 3.000000
57		57
58	```text	58	```text
59	annotated_args_refresh_add_custom	59	annotated_args_refresh_add_custom
60	-└── offline	60	+└── offline/cpp
61	├── CMakeLists.txt	61	├── CMakeLists.txt
62	├── run.sh	62	├── run.sh
63	├── ge	63	├── ge

Rexamples/custom_op/annotated_args_refresh_add_custom/offline/README_en.md→examples/custom_op/annotated_args_refresh_add_custom/offline/cpp/README_en.md+2-2

	@@ -33,7 +33,7 @@ Key implementation points:
33		33
34	## Quick Run	34	## Quick Run
35		35
36	-Run in `examples/custom_op/annotated_args_refresh_add_custom/offline`:	36	+Run in `examples/custom_op/annotated_args_refresh_add_custom/offline/cpp`:
37		37
38	```bash	38	```bash
39	bash run.sh	39	bash run.sh
	@@ -57,7 +57,7 @@ First element of output: 3.000000
57		57
58	```text	58	```text
59	annotated_args_refresh_add_custom	59	annotated_args_refresh_add_custom
60	-└── offline	60	+└── offline/cpp
61	├── CMakeLists.txt	61	├── CMakeLists.txt
62	├── run.sh	62	├── run.sh
63	├── ge	63	├── ge

Rexamples/custom_op/annotated_args_refresh_add_custom/offline/ge/add_custom_ir.h→examples/custom_op/annotated_args_refresh_add_custom/offline/cpp/ge/add_custom_ir.h+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/offline/ge/add_custom_kernel.cpp→examples/custom_op/annotated_args_refresh_add_custom/offline/cpp/ge/add_custom_kernel.cpp+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/offline/ge/custom_op.cpp→examples/custom_op/annotated_args_refresh_add_custom/offline/cpp/ge/custom_op.cpp+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/offline/ge/utils/compile_utils.cpp→examples/custom_op/annotated_args_refresh_add_custom/offline/cpp/ge/utils/compile_utils.cpp+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/offline/ge/utils/kernel_binary_map_utils.cpp→examples/custom_op/annotated_args_refresh_add_custom/offline/cpp/ge/utils/kernel_binary_map_utils.cpp+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/offline/graph_build/main.cc→examples/custom_op/annotated_args_refresh_add_custom/offline/cpp/graph_build/main.cc+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/offline/model_exec/main.cc→examples/custom_op/annotated_args_refresh_add_custom/offline/cpp/model_exec/main.cc+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/offline/run.sh→examples/custom_op/annotated_args_refresh_add_custom/offline/cpp/run.sh+0-0

文件重命名但无更改。

Aexamples/custom_op/annotated_args_refresh_add_custom/offline/python/CMakeLists.txt+59-0

@@ -0,0 +1,59 @@
	1	+cmake_minimum_required(VERSION 3.16)
	2	+project(annotated_args_refresh_add_offline_python LANGUAGES CXX)
	3	+
	4	+set(CMAKE_CXX_STANDARD 17)
	5	+set(CMAKE_CXX_STANDARD_REQUIRED ON)
	6	+set(CMAKE_CXX_EXTENSIONS OFF)
	7	+set(ES_OUTPUT_PATH "${CMAKE_BINARY_DIR}/es_output")
	8	+set(OPP_OUTPUT_PATH "${CMAKE_BINARY_DIR}/opp")
	9	+set(ASCEND_HOME_PATH "$ENV{ASCEND_HOME_PATH}")
	10	+if(NOT ASCEND_HOME_PATH)
	11	+ message(FATAL_ERROR "ASCEND_HOME_PATH is empty. Please source CANN set_env.sh first.")
	12	+endif()
	13	+
	14	+list(APPEND CMAKE_MODULE_PATH "${ASCEND_HOME_PATH}/include/ge/cmake")
	15	+find_package(GenerateEsPackage REQUIRED)
	16	+add_compile_definitions(_GLIBCXX_USE_CXX11_ABI=0 google=ascend_private)
	17	+
	18	+add_library(annotated_add_custom_op_proto SHARED proto/add_custom.cc)
	19	+target_compile_definitions(annotated_add_custom_op_proto PRIVATE OP_PROTO_LIB)
	20	+target_compile_options(annotated_add_custom_op_proto PRIVATE -fvisibility=hidden)
	21	+target_include_directories(annotated_add_custom_op_proto PRIVATE
	22	+ "${ASCEND_HOME_PATH}/include"
	23	+ "${ASCEND_HOME_PATH}/include/external"
	24	+)
	25	+target_link_libraries(annotated_add_custom_op_proto PRIVATE
	26	+ "${ASCEND_HOME_PATH}/lib64/libopp_registry.so"
	27	+)
	28	+
	29	+if(CMAKE_SYSTEM_NAME STREQUAL "Windows")
	30	+ set(OPP_OS_TYPE "windows")
	31	+else()
	32	+ set(OPP_OS_TYPE "linux")
	33	+endif()
	34	+string(TOLOWER "${CMAKE_SYSTEM_PROCESSOR}" PROCESSOR_LOWER)
	35	+if(PROCESSOR_LOWER MATCHES "^(aarch64\|arm64)$")
	36	+ set(OPP_CPU_TYPE "aarch64")
	37	+elseif(PROCESSOR_LOWER MATCHES "^(x86_64\|amd64)$")
	38	+ set(OPP_CPU_TYPE "x86_64")
	39	+else()
	40	+ set(OPP_CPU_TYPE "${PROCESSOR_LOWER}")
	41	+endif()
	42	+set(OPP_PROTO_OUTPUT_PATH "${OPP_OUTPUT_PATH}/op_proto/custom")
	43	+set(CUSTOM_OP_OUTPUT_PATH "${OPP_OUTPUT_PATH}/op_graph/lib/${OPP_OS_TYPE}/${OPP_CPU_TYPE}")
	44	+set_target_properties(annotated_add_custom_op_proto PROPERTIES
	45	+ OUTPUT_NAME cust_opapi
	46	+ LIBRARY_OUTPUT_DIRECTORY "${OPP_PROTO_OUTPUT_PATH}"
	47	+ RUNTIME_OUTPUT_DIRECTORY "${OPP_PROTO_OUTPUT_PATH}"
	48	+ ARCHIVE_OUTPUT_DIRECTORY "${OPP_PROTO_OUTPUT_PATH}"
	49	+)
	50	+add_custom_command(TARGET annotated_add_custom_op_proto POST_BUILD
	51	+ COMMAND ${CMAKE_COMMAND} -E make_directory "${CUSTOM_OP_OUTPUT_PATH}"
	52	+ COMMAND ${CMAKE_COMMAND} -E copy "$<TARGET_FILE:annotated_add_custom_op_proto>" "${CUSTOM_OP_OUTPUT_PATH}"
	53	+)
	54	+
	55	+add_es_library_and_whl(
	56	+ ES_LINKABLE_AND_ALL_TARGET es_custom
	57	+ OPP_PROTO_TARGET annotated_add_custom_op_proto
	58	+ OUTPUT_PATH ${ES_OUTPUT_PATH}
	59	+)

Aexamples/custom_op/annotated_args_refresh_add_custom/offline/python/README.md+47-0

@@ -0,0 +1,47 @@
	1	+# Python 离线自定义算子样例
	2	+
	3	+本样例演示使用 Python 装饰器完成 `AnnotatedAddCustom` 的原型注册、`infer_meta`、编译和 `declare_launch_args`，再生成 AIR/OM 并通过 ACL 在 NPU 上执行。`declare_launch_args` 只在 ATC 编译期运行，离线运行期只消费已经生成的任务描述。Python 回调复用 Ascend C kernel 源码并生成 device binary。
	4	+
	5	+## 目录
	6	+
	7	+```text
	8	+offline/python
	9	+├── CMakeLists.txt # 构建 custom OPP 注册库和 ES wheel
	10	+├── run.sh # kernel、ES、AIR、ATC 和 ACL 验证入口
	11	+├── proto/ # gen_esb 使用的 C++ 构图原型
	12	+└── src/
	13	+ ├── build_graph.py # Python 构图并生成 AIR
	14	+ ├── run_model.py # ACL 两轮离线 NPU 执行
	15	+ └── ge/annotated_add_custom.py # Python 原型和 infer_meta
	16	+```
	17	+
	18	+## 依赖与运行
	19	+
	20	+需要已安装并配置的 CANN、ATC、CMake、Python 3/pip、numpy 和可用 NPU。
	21	+
	22	+`proto/add_custom.h` 中的 C++ `REG_OP` 仅用于 `gen_esb` 生成 `ge.es.custom.AnnotatedAddCustom` 构图接口；运行时原型、`infer_meta`、`compile` 和 `declare_launch_args` 均由 `src/ge/annotated_add_custom.py` 的 Python 装饰器提供。
	23	+
	24	+```bash
	25	+source /path/to/cann/set_env.sh
	26	+cd examples/custom_op/annotated_args_refresh_add_custom/offline/python
	27	+bash run.sh
	28	+```
	29	+
	30	+可覆盖 `SOC_VERSION`（默认 `Ascend910B1`）和 `DEVICE_ID`（默认 `0`）。
	31	+
	32	+## 执行链路
	33	+
	34	+1. 构建 custom OPP 和 ES wheel，`build_graph.py` 生成 AIR。
	35	+2. ATC 导入 Python 模块执行 `infer_meta`、`compile` 和 `declare_launch_args`，将 binary 和地址布局写入
	36	+ `build/annotated_add.om`。
	37	+3. `run_model.py` 使用 ACL 加载 OM，分两轮创建独立数据集并执行，验证 AnnotatedArgs 地址刷新。
	38	+
	39	+运行日志中的 `NPU_TWO_ROUND_VALIDATION=PASS` 表示两轮输出均通过校验。ATC 日志应包含 Python 模块加载、
	40	+`infer_meta`、Python compile 和地址声明日志；OM 运行期不应再次导入 Python 模块。
	41	+
	42	+Add kernel 使用 Ascend C 编写，源码位于 `offline/cpp/ge/add_custom_kernel.cpp`。开发者可参考 CANN
	43	+Ascend C 文档编写 kernel，Python `compile` 回调负责生成并交付 binary。
	44	+
	45	+## callback 约束
	46	+
	47	+`append_input` 和 `append_output` 使用当前节点 input/output 的平铺 index。`AnnotatedArgsContext`、Tensor、workspace 和 args builder 都是 callback 期间的 borrowed 对象，不能逃逸；args builder 在 `add_launch` 后已 consumed，不能复用。

Aexamples/custom_op/annotated_args_refresh_add_custom/offline/python/README_en.md+55-0

@@ -0,0 +1,55 @@
	1	+# Python Offline Custom Operator
	2	+
	3	+This sample uses Python decorators to register the `AnnotatedAddCustom` prototype, `infer_meta`, `compile`, and
	4	+`declare_launch_args`. It generates AIR/OM and runs the model on an NPU through ACL. The declaration callback runs
	5	+only during ATC compilation; offline execution consumes the generated task description. The Python callback compiles
	6	+and serializes the Ascend C kernel binary.
	7	+
	8	+## Directory
	9	+
	10	+```text
	11	+offline/python
	12	+├── CMakeLists.txt # Build the custom OPP library and ES wheel
	13	+├── run.sh # Kernel, ES, AIR, ATC, and ACL validation entry point
	14	+├── proto/ # C++ graph prototype consumed by gen_esb
	15	+└── src/
	16	+ ├── build_graph.py # Build the Python graph and save AIR
	17	+ ├── run_model.py # Two-round ACL execution on an NPU
	18	+ └── ge/annotated_add_custom.py # Python prototype and infer_meta
	19	+```
	20	+
	21	+## Requirements and usage
	22	+
	23	+Install and configure CANN, ATC, CMake, Python 3/pip, numpy, and an available NPU.
	24	+
	25	+The C++ `REG_OP` in `proto/add_custom.h` is used only by `gen_esb` to generate the
	26	+`ge.es.custom.AnnotatedAddCustom` graph-building interface. The runtime prototype, `infer_meta`, `compile`, and
	27	+`declare_launch_args` are supplied by the Python decorators in `src/ge/annotated_add_custom.py`.
	28	+
	29	+```bash
	30	+source /path/to/cann/set_env.sh
	31	+cd examples/custom_op/annotated_args_refresh_add_custom/offline/python
	32	+bash run.sh
	33	+```
	34	+
	35	+The script accepts `SOC_VERSION` (default `Ascend910B1`) and `DEVICE_ID` (default `0`).
	36	+
	37	+## Execution flow
	38	+
	39	+1. Build the custom OPP and ES wheel; `build_graph.py` generates AIR.
	40	+2. ATC imports the Python module and invokes `infer_meta`, `compile`, and `declare_launch_args` to write the Ascend C
	41	+ binary and address layout into `build/annotated_add.om`.
	42	+3. `run_model.py` loads the OM with ACL, creates independent datasets for two rounds, and validates address refresh.
	43	+
	44	+`NPU_TWO_ROUND_VALIDATION=PASS` in the runtime log indicates that both rounds passed. The ATC log should contain
	45	+the Python module-load, `infer_meta`, compile, and address-declaration markers. These
	46	+compile-time markers must not appear during OM execution.
	47	+
	48	+The Add kernel is written with Ascend C in `offline/cpp/ge/add_custom_kernel.cpp`. Use the CANN Ascend C
	49	+documentation when developing kernels and ACL RTC when delivering their binaries.
	50	+
	51	+## Callback constraints
	52	+
	53	+Use the flattened input/output index of the current node with `append_input` and `append_output`.
	54	+`AnnotatedArgsContext`, tensors, workspace, and the argument builder are borrowed objects valid only during the
	55	+callback and must not escape. The argument builder is consumed by `add_launch` and cannot be reused.

Rexamples/custom_op/args_refresh_add_custom/python/proto/add_custom.cc→examples/custom_op/annotated_args_refresh_add_custom/offline/python/proto/add_custom.cc+3-1

	@@ -10,4 +10,6 @@
10		10
11	#include "add_custom.h"	11	#include "add_custom.h"
12		12
13	-namespace ge {}	13	+// Keep this translation unit as the OPP proto library input for gen_esb.
		14	+// Runtime prototype and infer-meta implementations are registered by the
		15	+// Python decorators in src/ge/annotated_add_custom.py.

Rexamples/custom_op/annotated_args_refresh_add_custom/python/proto/add_custom.h→examples/custom_op/annotated_args_refresh_add_custom/offline/python/proto/add_custom.h+1-2

	@@ -12,13 +12,12 @@
12	#define EXAMPLES_CUSTOM_OP_ANNOTATED_ARGS_REFRESH_ADD_CUSTOM_PYTHON_PROTO_ADD_CUSTOM_H_	12	#define EXAMPLES_CUSTOM_OP_ANNOTATED_ARGS_REFRESH_ADD_CUSTOM_PYTHON_PROTO_ADD_CUSTOM_H_
13		13
14	#include "graph/operator_reg.h"	14	#include "graph/operator_reg.h"
15	-#include "register/op_impl_registry.h"
16		15
17	namespace ge {	16	namespace ge {
18	REG_OP(AnnotatedAddCustom)	17	REG_OP(AnnotatedAddCustom)
19	.INPUT(x1, "T")	18	.INPUT(x1, "T")
20	.INPUT(x2, "T")	19	.INPUT(x2, "T")
21	- .OUTPUT(y, "T")	20	+ .OUTPUT(output0, "T")
22	.DATATYPE(T, TensorType({DT_FLOAT}))	21	.DATATYPE(T, TensorType({DT_FLOAT}))
23	.OP_END_FACTORY_REG(AnnotatedAddCustom);	22	.OP_END_FACTORY_REG(AnnotatedAddCustom);
24	} // namespace ge	23	} // namespace ge

Rexamples/custom_op/annotated_args_refresh_add_custom/python/run.sh→examples/custom_op/annotated_args_refresh_add_custom/offline/python/run.sh+40-103

	@@ -1,135 +1,72 @@
1	#!/usr/bin/env bash	1	#!/usr/bin/env bash
2	# -----------------------------------------------------------------------------------------------------------	2	# -----------------------------------------------------------------------------------------------------------
3	# Copyright (c) 2026 Huawei Technologies Co., Ltd.	3	# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4	-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of	4	+# This program is free software; you can redistribute it and/or modify it under the terms and conditions of
5	# CANN Open Software License Agreement Version 2.0 (the "License").	5	# 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.	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,	7	# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
8	# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.	8	# INCLUDING 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.	9	# See LICENSE in the root of the software repository for the full text of the License.
10	# -----------------------------------------------------------------------------------------------------------	10	# -----------------------------------------------------------------------------------------------------------
11	-
12	set -euo pipefail	11	set -euo pipefail
13		12
14	SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"	13	SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
15	BUILD_DIR="${SCRIPT_DIR}/build"	14	BUILD_DIR="${SCRIPT_DIR}/build"
16	-KERNEL_SOURCE="${SCRIPT_DIR}/../online/add_custom_kernel/add_custom.asc"	15	+KERNEL_SOURCE="${SCRIPT_DIR}/../cpp/ge/add_custom_kernel.cpp"
17	-ADD_CUSTOM_NPU_ARCH="${ADD_CUSTOM_NPU_ARCH:-2201}"	16	+KERNEL_ASC_SOURCE="${BUILD_DIR}/add_custom_kernel.asc"
18	SOC_VERSION="${SOC_VERSION:-Ascend910B1}"	17	SOC_VERSION="${SOC_VERSION:-Ascend910B1}"
19	-DEVICE_ID="${DEVICE_ID:-0}"
20	HOST_OS="$(uname -s \| tr '[:upper:]' '[:lower:]')"	18	HOST_OS="$(uname -s \| tr '[:upper:]' '[:lower:]')"
21	HOST_ARCH="$(uname -m \| tr '[:upper:]' '[:lower:]')"	19	HOST_ARCH="$(uname -m \| tr '[:upper:]' '[:lower:]')"
22	-ATC_BIN=""	20	+if [[ "${HOST_OS}" == mingw* \|\| "${HOST_OS}" == msys* \|\| "${HOST_OS}" == cygwin* ]]; then HOST_OS="windows"; else HOST_OS="linux"; fi
23	-	21	+case "${HOST_ARCH}" in arm64) HOST_ARCH="aarch64" ;; amd64) HOST_ARCH="x86_64" ;; esac
24	-if [[ "${HOST_OS}" == mingw* \|\| "${HOST_OS}" == msys* \|\| "${HOST_OS}" == cygwin* ]]; then	22	+info() { echo "[INFO] $*"; }
25	- HOST_OS="windows"	23	+error() { echo "[ERROR] $*" >&2; }
26	-else	24	+require_command() { command -v "$1" >/dev/null 2>&1 \|\| { error "Required command was not found: $1"; exit 1; }; }
27	- HOST_OS="linux"	25	+require_file() { [[ -s "$1" ]] \|\| { error "Required output was not generated: $1"; exit 1; }; }
28	-fi
29	-case "${HOST_ARCH}" in
30	- arm64)
31	- HOST_ARCH="aarch64"
32	- ;;
33	- amd64)
34	- HOST_ARCH="x86_64"
35	- ;;
36	-esac
37	-
38	-info() {
39	- echo "[INFO] $*"
40	-}
41	-
42	-error() {
43	- echo "[ERROR] $*" >&2
44	-}
45	-
46	-require_command() {
47	- if ! command -v "$1" >/dev/null 2>&1; then
48	- error "Required command was not found: $1"
49	- exit 1
50	- fi
51	-}
52	-
53	-require_file() {
54	- if [[ ! -s "$1" ]]; then
55	- error "Required output was not generated: $1"
56	- exit 1
57	- fi
58	-}
59		26
60	if [[ -z "${ASCEND_HOME_PATH:-}" \|\| ! -d "${ASCEND_HOME_PATH}" ]]; then	27	if [[ -z "${ASCEND_HOME_PATH:-}" \|\| ! -d "${ASCEND_HOME_PATH}" ]]; then
61	- error "ASCEND_HOME_PATH is empty or not a directory. Please source CANN set_env.sh first."	28	+ error "ASCEND_HOME_PATH is empty or not a directory. Please source CANN set_env.sh first."; exit 1
62	- exit 1
63	fi	29	fi
64	-if [[ ! -f "${KERNEL_SOURCE}" ]]; then	30	+for command_name in cmake python3 atc; do require_command "${command_name}"; done
65	- error "Kernel source was not found: ${KERNEL_SOURCE}"	31	+python3 -m pip --version >/dev/null 2>&1 \|\| { error "python3 -m pip is unavailable"; exit 1; }
66	- exit 1
67	-fi
68	-for command_name in bisheng llvm-objcopy cmake python3 atc; do
69	- require_command "${command_name}"
70	-done
71	-if ! python3 -m pip --version >/dev/null 2>&1; then
72	- error "python3 -m pip is unavailable"
73	- exit 1
74	-fi
75	-ATC_BIN="$(command -v atc)"
76	-
77	mkdir -p "${BUILD_DIR}"	32	mkdir -p "${BUILD_DIR}"
78		33
79	-info "Step 1/5: compile the Ascend C kernel"	34	+info "Step 1/4: build the custom OPP and Python ES wheel"
80	-bisheng -c "${KERNEL_SOURCE}" -o "${BUILD_DIR}/add_custom.host.o" --npu-arch="dav-${ADD_CUSTOM_NPU_ARCH}"
81	-llvm-objcopy -O binary --only-section=.aicore_binary "${BUILD_DIR}/add_custom.host.o" "${BUILD_DIR}/add_custom.o"
82	-require_file "${BUILD_DIR}/add_custom.o"
83	-
84	-info "Step 2/5: build the custom OPP and Python ES wheel"
85	cmake -S "${SCRIPT_DIR}" -B "${BUILD_DIR}" -DCMAKE_BUILD_TYPE=Release	35	cmake -S "${SCRIPT_DIR}" -B "${BUILD_DIR}" -DCMAKE_BUILD_TYPE=Release
86	cmake --build "${BUILD_DIR}" --target build_es_custom -j8	36	cmake --build "${BUILD_DIR}" --target build_es_custom -j8
87	CUSTOM_OP_LIBRARY="${BUILD_DIR}/opp/op_graph/lib/${HOST_OS}/${HOST_ARCH}/libcust_opapi.so"	37	CUSTOM_OP_LIBRARY="${BUILD_DIR}/opp/op_graph/lib/${HOST_OS}/${HOST_ARCH}/libcust_opapi.so"
88	-if [[ "${HOST_OS}" == "windows" ]]; then	38	+if [[ "${HOST_OS}" == "windows" ]]; then CUSTOM_OP_LIBRARY="${BUILD_DIR}/opp/op_graph/lib/${HOST_OS}/${HOST_ARCH}/cust_opapi.dll"; fi
89	- CUSTOM_OP_LIBRARY="${BUILD_DIR}/opp/op_graph/lib/${HOST_OS}/${HOST_ARCH}/cust_opapi.dll"
90	-fi
91	WHEEL_PATH="${BUILD_DIR}/es_output/whl/es_custom-1.0.0-py3-none-any.whl"	39	WHEEL_PATH="${BUILD_DIR}/es_output/whl/es_custom-1.0.0-py3-none-any.whl"
92	-require_file "${CUSTOM_OP_LIBRARY}"	40	+require_file "${CUSTOM_OP_LIBRARY}"; require_file "${WHEEL_PATH}"
93	-require_file "${WHEEL_PATH}"
94		41
95	-info "Step 3/5: install the generated Python ES wheel"	42	+info "Step 2/4: build AIR and compile it with ATC"
96	python3 -m pip install --force-reinstall --upgrade --target "${BUILD_DIR}/whl_package" "${WHEEL_PATH}"	43	python3 -m pip install --force-reinstall --upgrade --target "${BUILD_DIR}/whl_package" "${WHEEL_PATH}"
97	-export PYTHONPATH="${BUILD_DIR}/whl_package:${PYTHONPATH:-}"	44	+export PYTHONPATH="${BUILD_DIR}/whl_package:${SCRIPT_DIR}/src:${PYTHONPATH:-}"
98	export LD_LIBRARY_PATH="${BUILD_DIR}/es_output/lib64:${LD_LIBRARY_PATH:-}"	45	export LD_LIBRARY_PATH="${BUILD_DIR}/es_output/lib64:${LD_LIBRARY_PATH:-}"
99	export ASCEND_CUSTOM_OPP_PATH="${BUILD_DIR}/opp:${SCRIPT_DIR}/src/ge"	46	export ASCEND_CUSTOM_OPP_PATH="${BUILD_DIR}/opp:${SCRIPT_DIR}/src/ge"
100	-	47	+cp "${KERNEL_SOURCE}" "${KERNEL_ASC_SOURCE}"
101	-info "Step 4/5: build AIR and compile it with ATC"	48	+require_file "${KERNEL_ASC_SOURCE}"
102	-rm -f "${BUILD_DIR}/annotated_add.air" "${BUILD_DIR}/annotated_add.om" \	49	+export GE_PYTHON_CUSTOM_OP_SOURCE="${KERNEL_ASC_SOURCE}"
103	- "${BUILD_DIR}/annotated_add.json" "${BUILD_DIR}/ge_check_op.json"	50	+rm -f "${BUILD_DIR}/annotated_add.air" "${BUILD_DIR}/annotated_add.om" "${BUILD_DIR}/annotated_add.json" "${BUILD_DIR}/ge_check_op.json"
104	-python3 "${SCRIPT_DIR}/src/build_graph.py"	51	+python3 "${SCRIPT_DIR}/src/build_graph.py"; require_file "${BUILD_DIR}/annotated_add.air"
105	-require_file "${BUILD_DIR}/annotated_add.air"	52	+(cd "${BUILD_DIR}" && atc --model="${BUILD_DIR}/annotated_add.air" --framework=1 --output="${BUILD_DIR}/annotated_add" \
106	-(	53	+ --soc_version="${SOC_VERSION}" --host_env_os="${HOST_OS}" --host_env_cpu="${HOST_ARCH}" 2>&1 \| tee "${BUILD_DIR}/atc.log")
107	- cd "${BUILD_DIR}"
108	- "${ATC_BIN}" \
109	- --model="${BUILD_DIR}/annotated_add.air" \
110	- --framework=1 \
111	- --output="${BUILD_DIR}/annotated_add" \
112	- --soc_version="${SOC_VERSION}" 2>&1 \| tee "${BUILD_DIR}/atc.log"
113	-)
114	require_file "${BUILD_DIR}/annotated_add.om"	54	require_file "${BUILD_DIR}/annotated_add.om"
115	-if ! grep -Fq "PY_ANNOTATED_ARGS_MODULE_LOADED=1" "${BUILD_DIR}/atc.log" \|\| \	55	+for marker in PY_ANNOTATED_ARGS_MODULE_LOADED=1; do
116	- ! grep -Fq "PY_ANNOTATED_ARGS_CALLBACK_ENTER=1" "${BUILD_DIR}/atc.log"; then	56	+ grep -Fq "${marker}" "${BUILD_DIR}/atc.log" \|\| { error "ATC marker is missing: ${marker}"; exit 1; }
117	- error "ATC did not execute the Python annotated-args callback"	57	+done
118	- exit 1	58	+for marker in "PY_COMPILE_CALLBACK_ENTER=1"; do
119	-fi	59	+ grep -Fq "${marker}" "${BUILD_DIR}/atc.log" \|\| { error "Python compile marker is missing: ${marker}"; exit 1; }
120	-
121	-info "Step 5/5: execute two address-refresh rounds on NPU"
122	-unset ASCEND_CUSTOM_OPP_PATH
123	-DEVICE_ID="${DEVICE_ID}" python3 "${SCRIPT_DIR}/src/run_model.py" 2>&1 \| tee "${BUILD_DIR}/runtime.log"
124	-for marker in ROUND_1_FIRST=3 ROUND_2_FIRST=9 NPU_TWO_ROUND_VALIDATION=PASS; do
125	- if ! grep -Fq "${marker}" "${BUILD_DIR}/runtime.log"; then
126	- error "Runtime validation marker is missing: ${marker}"
127	- exit 1
128	- fi
129	done	60	done
130	-if grep -Eq "PY_ANNOTATED_ARGS_(MODULE_LOADED\|CALLBACK_ENTER)=1" "${BUILD_DIR}/runtime.log"; then
131	- error "Runtime log contains compile-time Python callback markers"
132	- exit 1
133	-fi
134		61
135	-info "Annotated args Python ATC/NPU pipeline PASS"	62	+info "Step 3/4: execute two offline AnnotatedArgs address-refresh rounds on NPU"
		63	+unset ASCEND_CUSTOM_OPP_PATH
		64	+DEVICE_ID="${DEVICE_ID:-0}" python3 "${SCRIPT_DIR}/src/run_model.py" 2>&1 \| tee "${BUILD_DIR}/runtime.log"
		65	+for marker in ROUND_1_FIRST=3 ROUND_2_FIRST=9 NPU_TWO_ROUND_VALIDATION=PASS; do
		66	+ grep -Fq "${marker}" "${BUILD_DIR}/runtime.log" \|\| { error "Runtime validation marker is missing: ${marker}"; exit 1; }
		67	+done
		68	+if grep -Eq "PY_ANNOTATED_ARGS_(MODULE_LOADED\|INFER_META\|COMPILE_(ENTER\|EXIT)\|CALLBACK_ENTER)=1" \
		69	+ "${BUILD_DIR}/runtime.log"; then
		70	+ error "Runtime log contains compile-time Python callback markers"; exit 1
		71	+fi
		72	+info "Step 4/4: offline Python AIR/OM pipeline PASS"

Rexamples/custom_op/annotated_args_refresh_add_custom/python/src/build_graph.py→examples/custom_op/annotated_args_refresh_add_custom/offline/python/src/build_graph.py+2-5

	@@ -42,11 +42,8 @@ def build_graph():
42	format=Format.FORMAT_ND,	42	format=Format.FORMAT_ND,
43	shape=[NUM_ELEMENTS],	43	shape=[NUM_ELEMENTS],
44	)	44	)
45	- y = AnnotatedAddCustom(x1, x2)	45	+ output0 = AnnotatedAddCustom(x1, x2)
46	- y.set_shape([NUM_ELEMENTS]).set_format(Format.FORMAT_ND).set_data_type(	46	+ builder.set_graph_output(output0, 0)
47	- DataType.DT_FLOAT
48	- )
49	- builder.set_graph_output(y, 0)
50	return builder.build_and_reset()	47	return builder.build_and_reset()
51		48
52		49

Aexamples/custom_op/annotated_args_refresh_add_custom/offline/python/src/ge/annotated_add_custom.py+177-0

@@ -0,0 +1,177 @@
	1	+#!/usr/bin/env python3
	2	+# -- coding: utf-8 --
	3	+# -----------------------------------------------------------------------------------------------------------
	4	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	5	+# This program is free software; you can redistribute it and/or modify it under the terms and conditions of
	6	+# CANN Open Software License Agreement Version 2.0 (the "License").
	7	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	8	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	9	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	10	+# See LICENSE in the root of the software repository for the full text of the License.
	11	+# -----------------------------------------------------------------------------------------------------------
	12	+
	13	+"""Python prototype and AnnotatedArgs implementation for the offline sample."""
	14	+
	15	+import hashlib
	16	+import logging
	17	+import os
	18	+import subprocess
	19	+import tempfile
	20	+import threading
	21	+from pathlib import Path
	22	+from dataclasses import dataclass
	23	+
	24	+from ge.custom_op import (
	25	+ AnnotatedKernelLaunchInfo,
	26	+ get_compile_platform_info,
	27	+ get_declare_launch_args_ctx,
	28	+ register_op,
	29	+ register_op_impl,
	30	+)
	31	+from ge.runtime import Tensor, TensorDesc
	32	+
	33	+logging.basicConfig(level=logging.INFO, format="%(message)s")
	34	+_LOGGER = logging.getLogger(__name__)
	35	+_LOGGER.info("PY_ANNOTATED_ARGS_MODULE_LOADED=1")
	36	+
	37	+
	38	+@register_op(op_type="AnnotatedAddCustom")
	39	+def annotated_add_custom_infer_meta(x1: TensorDesc, x2: TensorDesc) -> TensorDesc:
	40	+ """Infer the output metadata for elementwise addition."""
	41	+
	42	+ _LOGGER.info("PY_ANNOTATED_ARGS_INFER_META_ENTER=1")
	43	+ if x1.shape.origin_shape.dims != x2.shape.origin_shape.dims:
	duhua25 天前可以加一下日志打印，作为example正常运行的凭证
	44	+ raise ValueError("AnnotatedAddCustom inputs must have the same shape")
	45	+ if x1.data_type != x2.data_type:
	46	+ raise ValueError("AnnotatedAddCustom inputs must have the same data type")
	47	+ output = TensorDesc(x1.shape, x1.data_type)
	48	+ _LOGGER.info("PY_ANNOTATED_ARGS_INFER_META_EXIT=1")
	49	+ return output
	50	+
	51	+
	52	+_KERNEL_NAME = "add_custom"
	53	+_BLOCK_SIZE = 1024
	54	+_KERNEL_SOURCE = Path(
	55	+ os.environ.get(
	56	+ "GE_PYTHON_CUSTOM_OP_SOURCE",
	57	+ str(
	58	+ Path(__file__).resolve().parents[3] / "cpp" / "ge" / "add_custom_kernel.cpp"
	59	+ ),
	60	+ )
	61	+)
	62	+
	63	+
	64	+@dataclass(frozen=True)
	65	+class _Artifact:
	66	+ binary: bytes
	67	+ block_dim: int
	68	+
	69	+
	70	+def _dims(tensor: Tensor):
	71	+ return tuple(int(dim) for dim in tensor.storage_shape.dims)
	72	+
	73	+
	74	+def _validate(x: Tensor, y: Tensor, z: Tensor) -> None:
	75	+ if _dims(x) != _dims(y) or _dims(x) != _dims(z):
	76	+ raise ValueError("AnnotatedAddCustom requires matching tensor shapes")
	77	+ if x.data_type != y.data_type or x.data_type != z.data_type:
	78	+ raise ValueError("AnnotatedAddCustom requires matching data types")
	79	+ if not _dims(x) or any(dim <= 0 for dim in _dims(x)):
	80	+ raise ValueError("AnnotatedAddCustom requires a concrete positive shape")
	81	+ elements = 1
	82	+ for dim in _dims(x):
	83	+ elements *= dim
	84	+ if elements % _BLOCK_SIZE != 0:
	85	+ raise ValueError(
	86	+ "the sample kernel requires an element count divisible by 1024"
	87	+ )
	88	+
	89	+
	90	+def _arch(platform):
	91	+ value = str(platform.get_platform_resource("version", "NpuArch")).strip()
	92	+ return value[4:] if value.startswith("dav-") else value
	93	+
	94	+
	95	+def _compile_kernel(platform) -> bytes:
	96	+ ascend_home = os.environ.get("ASCEND_HOME_PATH")
	97	+ if not ascend_home or not _KERNEL_SOURCE.is_file():
	98	+ raise RuntimeError(
	99	+ "ASCEND_HOME_PATH and the Ascend C kernel source are required"
	100	+ )
	101	+ include = Path(ascend_home) / "asc" / "include"
	102	+ out = Path(os.environ.get("GE_PYTHON_CUSTOM_OP_BUILD_DIR", tempfile.gettempdir()))
	103	+ out.mkdir(parents=True, exist_ok=True)
	104	+ digest = hashlib.sha256(_KERNEL_SOURCE.read_bytes()).hexdigest()[:16]
	105	+ host = out / ("annotated_add_custom_" + digest + ".o")
	106	+ binary = out / ("annotated_add_custom_" + digest + ".aicore.o")
	107	+ if not binary.exists():
	108	+ subprocess.run(
	109	+ [
	110	+ "bisheng",
	111	+ "-c",
	112	+ str(_KERNEL_SOURCE),
	113	+ "-o",
	114	+ str(host),
	115	+ "--npu-arch=dav-" + _arch(platform),
	116	+ "-I" + str(include),
	117	+ ],
	118	+ check=True,
	119	+ )
	120	+ subprocess.run(
	121	+ [
	122	+ "llvm-objcopy",
	123	+ "-O",
	124	+ "binary",
	125	+ "--only-section=.aicore_binary",
	126	+ str(host),
	127	+ str(binary),
	128	+ ],
	129	+ check=True,
	130	+ )
	131	+ data = binary.read_bytes()
	132	+ if not data:
	133	+ raise RuntimeError("Ascend C kernel binary is empty")
	134	+ return bytes(data)
	135	+
	136	+
	137	+@register_op_impl(op_type="AnnotatedAddCustom")
	138	+class AnnotatedAddCustom:
	139	+ """Compile the Ascend C kernel and publish declarative launch arguments."""
	140	+
	141	+ def __init__(self):
	142	+ self._artifacts = {}
	143	+ self._lock = threading.RLock()
	144	+
	145	+ def compile(self, x: Tensor, y: Tensor, z: Tensor) -> None:
	146	+ _validate(x, y, z)
	147	+ key = (_dims(x), str(x.data_type))
	148	+ with self._lock:
	149	+ if key not in self._artifacts:
	150	+ print("PY_COMPILE_CALLBACK_ENTER=1", flush=True)
	151	+ platform = get_compile_platform_info()
	152	+ elements = 1
	153	+ for dim in _dims(x):
	154	+ elements *= dim
	155	+ self._artifacts[key] = _Artifact(
	156	+ _compile_kernel(platform), elements // _BLOCK_SIZE
	157	+ )
	158	+
	159	+ def declare_launch_args(self, x: Tensor, y: Tensor, z: Tensor) -> None:
	160	+ _validate(x, y, z)
	161	+ artifact = self._artifacts.get((_dims(x), str(x.data_type)))
	162	+ if artifact is None:
	163	+ raise RuntimeError("AnnotatedAddCustom compile cache miss")
	164	+ ctx = get_declare_launch_args_ctx()
	165	+ args = ctx.create_kernel_args()
	166	+ args.append_input(0, x)
	167	+ args.append_input(1, y)
	168	+ args.append_output(0, z)
	169	+ ctx.add_launch(
	170	+ AnnotatedKernelLaunchInfo(
	171	+ kernel_name=_KERNEL_NAME,
	172	+ kernel_bin=artifact.binary,
	173	+ block_dim=artifact.block_dim,
	174	+ stream_id=ctx.get_stream_id(),
	175	+ ),
	176	+ args,
	177	+ )

Rexamples/custom_op/annotated_args_refresh_add_custom/python/src/run_model.py→examples/custom_op/annotated_args_refresh_add_custom/offline/python/src/run_model.py+3-3

	@@ -175,9 +175,9 @@ def create_round(model_desc: Any, x_value: float, y_value: float) -> RoundData:
175	"""Allocate one complete, non-reused input/output dataset pair."""	175	"""Allocate one complete, non-reused input/output dataset pair."""
176	import acl	176	import acl
177		177
178	- input_sizes = [	178	+ input_sizes = []
179	- acl.mdl.get_input_size_by_index(model_desc, index) for index in range(2)	179	+ for index in range(2):
180	- ]	180	+ input_sizes.append(acl.mdl.get_input_size_by_index(model_desc, index))
181	output_size = acl.mdl.get_output_size_by_index(model_desc, 0)	181	output_size = acl.mdl.get_output_size_by_index(model_desc, 0)
182	inputs = create_dataset(input_sizes, [x_value, y_value])	182	inputs = create_dataset(input_sizes, [x_value, y_value])
183	try:	183	try:

Rexamples/custom_op/annotated_args_refresh_add_custom/online/CMakeLists.txt→examples/custom_op/annotated_args_refresh_add_custom/online/cpp/CMakeLists.txt+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/online/README.md→examples/custom_op/annotated_args_refresh_add_custom/online/cpp/README.md+3-3

	@@ -26,7 +26,7 @@
26		26
27	- 已正确安装并配置 CANN 环境，例如执行过 `source ${ASCEND_HOME_PATH}/set_env.sh`。	27	- 已正确安装并配置 CANN 环境，例如执行过 `source ${ASCEND_HOME_PATH}/set_env.sh`。
28	- 当前环境具备 `ACL`、`GE`、`Graph` 相关头文件与库。	28	- 当前环境具备 `ACL`、`GE`、`Graph` 相关头文件与库。
29	-- 参考 [安装指导](../../../../docs/zh/quick_install.md) 完成 toolkit 和 ops 包安装。	29	+- 参考 [安装指导](../../../../../docs/zh/quick_install.md) 完成 toolkit 和 ops 包安装。
30		30
31	### 框架与插件	31	### 框架与插件
32		32
	@@ -45,7 +45,7 @@
45		45
46	## 快速运行	46	## 快速运行
47		47
48	-在 `examples/custom_op/annotated_args_refresh_add_custom/online` 目录下执行：	48	+在 `examples/custom_op/annotated_args_refresh_add_custom/online/cpp` 目录下执行：
49		49
50	### 推荐方式	50	### 推荐方式
51		51
	@@ -89,7 +89,7 @@ cd ..
89		89
90	```text	90	```text
91	annotated_args_refresh_add_custom	91	annotated_args_refresh_add_custom
92	-└── online	92	+└── online/cpp
93	├── CMakeLists.txt	93	├── CMakeLists.txt
94	├── README.md	94	├── README.md
95	├── README_en.md	95	├── README_en.md

Rexamples/custom_op/annotated_args_refresh_add_custom/online/README_en.md→examples/custom_op/annotated_args_refresh_add_custom/online/cpp/README_en.md+3-3

	@@ -26,7 +26,7 @@ This sample defines two functionally identical Add custom operators. Both accept
26		26
27	- The CANN environment is installed and configured, for example by running `source ${ASCEND_HOME_PATH}/set_env.sh`.	27	- The CANN environment is installed and configured, for example by running `source ${ASCEND_HOME_PATH}/set_env.sh`.
28	- The environment provides the required `ACL`, `GE`, and `Graph` headers and libraries.	28	- The environment provides the required `ACL`, `GE`, and `Graph` headers and libraries.
29	-- Refer to the [Installation Guide](../../../../docs/en/quick_install.md) to install the toolkit and ops packages.	29	+- Refer to the [Installation Guide](../../../../../docs/en/quick_install.md) to install the toolkit and ops packages.
30		30
31	### Frameworks and Plugins	31	### Frameworks and Plugins
32		32
	@@ -45,7 +45,7 @@ This sample defines two functionally identical Add custom operators. Both accept
45		45
46	## Quick Run	46	## Quick Run
47		47
48	-Run the following commands in `examples/custom_op/annotated_args_refresh_add_custom/online`:	48	+Run the following commands in `examples/custom_op/annotated_args_refresh_add_custom/online/cpp`:
49		49
50	### Recommended Method	50	### Recommended Method
51		51
	@@ -89,7 +89,7 @@ cd ..
89		89
90	```text	90	```text
91	annotated_args_refresh_add_custom	91	annotated_args_refresh_add_custom
92	-└── online	92	+└── online/cpp
93	├── CMakeLists.txt	93	├── CMakeLists.txt
94	├── README.md	94	├── README.md
95	├── README_en.md	95	├── README_en.md

Rexamples/custom_op/annotated_args_refresh_add_custom/online/add_custom_kernel/add_custom.asc→examples/custom_op/annotated_args_refresh_add_custom/online/cpp/add_custom_kernel/add_custom.asc+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/online/add_custom_kernel/add_custom_kernel.h→examples/custom_op/annotated_args_refresh_add_custom/online/cpp/add_custom_kernel/add_custom_kernel.h+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/online/ge/add_custom_ir.h→examples/custom_op/annotated_args_refresh_add_custom/online/cpp/ge/add_custom_ir.h+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/online/ge/custom_op.cpp→examples/custom_op/annotated_args_refresh_add_custom/online/cpp/ge/custom_op.cpp+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/online/ge/utils/log.h→examples/custom_op/annotated_args_refresh_add_custom/online/cpp/ge/utils/log.h+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/online/ge/utils/rtc_kernel_loader.cpp→examples/custom_op/annotated_args_refresh_add_custom/online/cpp/ge/utils/rtc_kernel_loader.cpp+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/online/ge/utils/rtc_kernel_loader.h→examples/custom_op/annotated_args_refresh_add_custom/online/cpp/ge/utils/rtc_kernel_loader.h+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/online/run.sh→examples/custom_op/annotated_args_refresh_add_custom/online/cpp/run.sh+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/online/session_run/main.cc→examples/custom_op/annotated_args_refresh_add_custom/online/cpp/session_run/main.cc+0-0

文件重命名但无更改。

Rexamples/custom_op/annotated_args_refresh_add_custom/python/CMakeLists.txt→examples/custom_op/annotated_args_refresh_add_custom/online/python/CMakeLists.txt+1-1

	@@ -1,5 +1,5 @@
1	cmake_minimum_required(VERSION 3.16)	1	cmake_minimum_required(VERSION 3.16)
2	-project(annotated_args_refresh_add_python LANGUAGES CXX)	2	+project(annotated_args_refresh_add_online_python LANGUAGES CXX)
3		3
4	set(CMAKE_CXX_STANDARD 17)	4	set(CMAKE_CXX_STANDARD 17)
5	set(CMAKE_CXX_STANDARD_REQUIRED ON)	5	set(CMAKE_CXX_STANDARD_REQUIRED ON)

Aexamples/custom_op/annotated_args_refresh_add_custom/online/python/README.md+55-0

@@ -0,0 +1,55 @@
	1	+# Python 在线自定义算子性能对比样例
	2	+
	3	+本样例使用 Python 装饰器注册 `AnnotatedAddCustom` 和 `NoRefreshAddCustom` 的原型与 `infer_meta`，
	4	+在两张在线 GE 图中对比声明式地址刷新和不声明地址刷新的执行性能。两条链路使用同一份 Ascend C Add kernel，
	5	+脚本以两组独立的 device Tensor 交替运行，完成精度校验、预热和计时，不生成 AIR/OM。
	6	+
	7	+## 目录
	8	+
	9	+```text
	10	+online/python
	11	+├── CMakeLists.txt # 构建 custom OPP 注册库和 ES wheel
	12	+├── run.sh # 构建并执行在线性能对比
	13	+├── proto/ # gen_esb 使用的 C++ 构图原型
	14	+└── src/
	15	+ ├── run.py # 两张在线 GE 图的精度和性能对比
	16	+ └── ge/
	17	+ └── annotated_add_custom.py # Python 原型和 infer_meta
	18	+```
	19	+
	20	+## 依赖与运行
	21	+
	22	+需要已安装并配置的 CANN、CMake、Python 3/pip 和可用 NPU。
	23	+
	24	+`proto/add_custom.h` 中的 C++ `REG_OP` 仅用于 `gen_esb` 生成构图接口；运行时原型和 `infer_meta`
	25	+由 Python 装饰器提供；`AnnotatedAddCustom` 的 `compile`、`declare_launch_args` 和 `NoRefreshAddCustom` 的
	26	+`execute` 均由 Python `register_op_impl` 提供。两个实现复用 `cpp/add_custom_kernel/add_custom.asc` 生成的
	27	+Ascend C binary。
	28	+
	29	+本目录不编译 `online/cpp/ge/custom_op.cpp`；该文件仅作为 C++ 对照样例，Python 样例的算子执行全部由 Python 回调完成。
	30	+
	31	+```bash
	32	+source /path/to/cann/set_env.sh
	33	+cd examples/custom_op/annotated_args_refresh_add_custom/online/python
	34	+bash run.sh
	35	+```
	36	+
	37	+可通过 `DEVICE_ID`（默认 `0`）选择 NPU：
	38	+
	39	+```bash
	40	+DEVICE_ID=1 bash run.sh
	41	+```
	42	+
	43	+## 执行链路
	44	+
	45	+1. 构建 custom OPP 和 ES wheel，生成两个构图接口。
	46	+2. Python `compile` 回调在 GE 编译图时编译并保存 Ascend C Add kernel binary。
	47	+3. Python `declare_launch_args` 回调声明输入、输出地址槽和 kernel launch；GE 在重复执行时刷新地址。
	48	+4. Python `execute` 为 schema 输出申请 Tensor，并通过 ACL Python API 下发同一 Ascend C kernel，作为无地址声明基线。
	49	+5. `run.py` 为两张图交替传入两组 device Tensor，分别完成精度校验、5 次预热和 100 次计时。
	50	+
	51	+日志会输出 `AnnotatedAddCustom`、`NoRefreshAddCustom` 的总耗时、平均耗时和 `Annotated speedup`；
	52	+`NPU_EXECUTION=PASS` 表示两条链路均通过精度和执行校验。
	53	+
	54	+Add kernel 使用 Ascend C 编写，源码位于 `online/cpp/add_custom_kernel/add_custom.asc`。开发者可参考
	55	+CANN Ascend C 文档编写并通过 ACL RTC 编译自定义 kernel。

Aexamples/custom_op/annotated_args_refresh_add_custom/online/python/README_en.md+58-0

@@ -0,0 +1,58 @@
	1	+# Python Online Custom Operator Performance Comparison
	2	+
	3	+This sample uses Python decorators to register the prototypes and `infer_meta` functions of `AnnotatedAddCustom` and `NoRefreshAddCustom`. Two online GE
	4	+graphs compare declarative address refresh with a path that does not declare
	5	+refreshable addresses. Both paths use the same Ascend C Add kernel. The script validates results and measures
	6	+execution without generating AIR/OM.
	7	+
	8	+## Directory
	9	+
	10	+```text
	11	+online/python
	12	+├── CMakeLists.txt # Build the custom OPP library and ES wheel
	13	+├── run.sh # Build and run the online comparison
	14	+├── proto/ # C++ graph prototypes consumed by gen_esb
	15	+└── src/
	16	+ ├── run.py # Validate and benchmark the two online graphs
	17	+ └── ge/annotated_add_custom.py # Python prototypes and infer_meta
	18	+```
	19	+
	20	+## Requirements and usage
	21	+
	22	+Install and configure CANN, CMake, Python 3/pip, and an available NPU.
	23	+
	24	+The C++ `REG_OP` declarations in `proto/add_custom.h` are used only by `gen_esb` to generate the graph APIs.
	25	+The runtime prototypes and `infer_meta` functions of both operators are supplied by Python decorators. Python `register_op_impl`
	26	+supplies `compile`, `declare_launch_args`, and `execute` for both operators. They reuse the Ascend C binary generated
	27	+from `cpp/add_custom_kernel/add_custom.asc`.
	28	+
	29	+This directory does not compile `online/cpp/ge/custom_op.cpp`; that file is provided only as the C++ comparison sample.
	30	+All operator execution in this Python sample is implemented by Python callbacks.
	31	+
	32	+```bash
	33	+source /path/to/cann/set_env.sh
	34	+cd examples/custom_op/annotated_args_refresh_add_custom/online/python
	35	+bash run.sh
	36	+```
	37	+
	38	+Use `DEVICE_ID` (default `0`) to select the NPU:
	39	+
	40	+```bash
	41	+DEVICE_ID=1 bash run.sh
	42	+```
	43	+
	44	+## Execution flow
	45	+
	46	+1. The custom OPP and ES wheel are built to generate both Python graph APIs.
	47	+2. The Python `compile` callback compiles and retains the Ascend C Add kernel binary while GE compiles the graph.
	48	+3. The Python `declare_launch_args` callback declares the input/output address slots and kernel launch, allowing GE to refresh addresses during
	49	+ repeated execution.
	50	+4. Python `execute` allocates the schema output and launches the same Ascend C kernel through the ACL Python API as the
	51	+ no-declaration baseline.
	52	+5. `run.py` validates both graphs, performs five warm-up iterations, and measures 100 iterations.
	53	+
	54	+The log reports total and average times for both operators plus `Annotated speedup`.
	55	+`NPU_EXECUTION=PASS` indicates that both paths passed validation and execution.
	56	+
	57	+The Add kernel is written with Ascend C in `online/cpp/add_custom_kernel/add_custom.asc`. Use the CANN Ascend C
	58	+documentation when developing kernels for the target hardware.

Rexamples/custom_op/annotated_args_refresh_add_custom/python/proto/add_custom.cc→examples/custom_op/annotated_args_refresh_add_custom/online/python/proto/add_custom.cc+4-18

	@@ -1,6 +1,6 @@
1	/**	1	/**
2	* Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	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.	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,	6	* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	@@ -10,20 +10,6 @@
10		10
11	#include "add_custom.h"	11	#include "add_custom.h"
12		12
13	-namespace {	13	+// Keep this translation unit as the OPP proto library input for gen_esb.
14	-ge::graphStatus InferShape(gert::InferShapeContext *ctx) {	14	+// Runtime prototype and infer-meta implementations are registered by the
15	- const auto *input_shape = ctx->GetInputShape(0U);	15	+// Python decorators in src/ge/annotated_add_custom.py.
16	- auto *output_shape = ctx->GetOutputShape(0U);
17	- if ((input_shape == nullptr) \|\| (output_shape == nullptr)) {
18	- return ge::GRAPH_FAILED;
19	- }
20	- output_shape = input_shape;
21	- return ge::GRAPH_SUCCESS;
22	-}
23	-
24	-ge::graphStatus InferDataType(gert::InferDataTypeContext *ctx) {
25	- return ctx->SetOutputDataType(0U, ctx->GetInputDataType(0U));
26	-}
27	-
28	-IMPL_OP(AnnotatedAddCustom).InferShape(InferShape).InferDataType(InferDataType);
29	-} // namespace

Rexamples/custom_op/args_refresh_add_custom/python/proto/add_custom.h→examples/custom_op/annotated_args_refresh_add_custom/online/python/proto/add_custom.h+16-10

	@@ -1,6 +1,6 @@
1	/**	1	/**
2	* Copyright (c) 2026 Huawei Technologies Co., Ltd.	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	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").	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.	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,	6	* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	@@ -8,19 +8,25 @@
8	* See LICENSE in the root of the software repository for the full text of the License.	8	* See LICENSE in the root of the software repository for the full text of the License.
9	*/	9	*/
10		10
11	-#ifndef EXAMPLES_CUSTOM_OP_ARGS_REFRESH_ADD_CUSTOM_PYTHON_PROTO_ADD_CUSTOM_H_	11	+#ifndef EXAMPLES_CUSTOM_OP_ANNOTATED_ARGS_REFRESH_ADD_CUSTOM_ONLINE_PYTHON_PROTO_ADD_CUSTOM_H_
12	-#define EXAMPLES_CUSTOM_OP_ARGS_REFRESH_ADD_CUSTOM_PYTHON_PROTO_ADD_CUSTOM_H_	12	+#define EXAMPLES_CUSTOM_OP_ANNOTATED_ARGS_REFRESH_ADD_CUSTOM_ONLINE_PYTHON_PROTO_ADD_CUSTOM_H_
13		13
14	#include "graph/operator_reg.h"	14	#include "graph/operator_reg.h"
15	-#include "register/op_impl_registry.h"
16		15
17	namespace ge {	16	namespace ge {
18	-REG_OP(AddPythonCustomOp)	17	+REG_OP(AnnotatedAddCustom)
19	- .INPUT(x, "T")	18	+ .INPUT(x1, "T")
20	- .INPUT(y, "T")	19	+ .INPUT(x2, "T")
21	- .OUTPUT(z, "T")	20	+ .OUTPUT(output0, "T")
22	.DATATYPE(T, TensorType({DT_FLOAT}))	21	.DATATYPE(T, TensorType({DT_FLOAT}))
23	- .OP_END_FACTORY_REG(AddPythonCustomOp);	22	+ .OP_END_FACTORY_REG(AnnotatedAddCustom);
		23	+
		24	+REG_OP(NoRefreshAddCustom)
		25	+ .INPUT(x1, "T")
		26	+ .INPUT(x2, "T")
		27	+ .OUTPUT(output0, "T")
		28	+ .DATATYPE(T, TensorType({DT_FLOAT}))
		29	+ .OP_END_FACTORY_REG(NoRefreshAddCustom);
24	} // namespace ge	30	} // namespace ge
25		31
26	-#endif // EXAMPLES_CUSTOM_OP_ARGS_REFRESH_ADD_CUSTOM_PYTHON_PROTO_ADD_CUSTOM_H_	32	+#endif // EXAMPLES_CUSTOM_OP_ANNOTATED_ARGS_REFRESH_ADD_CUSTOM_ONLINE_PYTHON_PROTO_ADD_CUSTOM_H_

Aexamples/custom_op/annotated_args_refresh_add_custom/online/python/run.sh+65-0

@@ -0,0 +1,65 @@
	1	+#!/usr/bin/env bash
	2	+# -----------------------------------------------------------------------------------------------------------
	3	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	4	+# This program is free software; you can redistribute it and/or modify it under the terms and conditions of
	5	+# 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,
	8	+# INCLUDING 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	+set -euo pipefail
	12	+
	13	+SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
	14	+BUILD_DIR="${SCRIPT_DIR}/build"
	15	+KERNEL_SOURCE="${SCRIPT_DIR}/../cpp/add_custom_kernel/add_custom.asc"
	16	+KERNEL_HOST_OBJECT="${BUILD_DIR}/add_custom.host.o"
	17	+KERNEL_BINARY="${BUILD_DIR}/add_custom.aicore.o"
	18	+NPU_ARCH="${ADD_CUSTOM_NPU_ARCH:-2201}"
	19	+HOST_OS="$(uname -s \| tr '[:upper:]' '[:lower:]')"
	20	+HOST_ARCH="$(uname -m \| tr '[:upper:]' '[:lower:]')"
	21	+if [[ "${HOST_OS}" == mingw* \|\| "${HOST_OS}" == msys* \|\| "${HOST_OS}" == cygwin* ]]; then HOST_OS="windows"; else HOST_OS="linux"; fi
	22	+case "${HOST_ARCH}" in arm64) HOST_ARCH="aarch64" ;; amd64) HOST_ARCH="x86_64" ;; esac
	23	+info() { echo "[INFO] $*"; }
	24	+error() { echo "[ERROR] $*" >&2; }
	25	+require_command() { command -v "$1" >/dev/null 2>&1 \|\| { error "Required command was not found: $1"; exit 1; }; }
	26	+require_file() { [[ -s "$1" ]] \|\| { error "Required output was not generated: $1"; exit 1; }; }
	27	+
	28	+if [[ -z "${ASCEND_HOME_PATH:-}" \|\| ! -d "${ASCEND_HOME_PATH}" ]]; then
	29	+ error "ASCEND_HOME_PATH is empty or not a directory. Please source CANN set_env.sh first."; exit 1
	30	+fi
	31	+for command_name in cmake python3 bisheng llvm-objcopy; do require_command "${command_name}"; done
	32	+python3 -m pip --version >/dev/null 2>&1 \|\| { error "python3 -m pip is unavailable"; exit 1; }
	33	+mkdir -p "${BUILD_DIR}"
	34	+
	35	+info "Step 1/4: compile the Ascend C kernel"
	36	+require_file "${KERNEL_SOURCE}"
	37	+mkdir -p "${BUILD_DIR}"
	38	+bisheng -c "${KERNEL_SOURCE}" -o "${KERNEL_HOST_OBJECT}" --npu-arch="dav-${NPU_ARCH}"
	39	+llvm-objcopy -O binary --only-section=.aicore_binary "${KERNEL_HOST_OBJECT}" "${KERNEL_BINARY}"
	40	+require_file "${KERNEL_BINARY}"
	41	+
	42	+info "Step 2/4: build the custom OPP and Python ES wheel"
	43	+cmake -S "${SCRIPT_DIR}" -B "${BUILD_DIR}" -DCMAKE_BUILD_TYPE=Release
	44	+cmake --build "${BUILD_DIR}" --target build_es_custom -j8
	45	+CUSTOM_OP_LIBRARY="${BUILD_DIR}/opp/op_graph/lib/${HOST_OS}/${HOST_ARCH}/libcust_opapi.so"
	46	+if [[ "${HOST_OS}" == "windows" ]]; then CUSTOM_OP_LIBRARY="${BUILD_DIR}/opp/op_graph/lib/${HOST_OS}/${HOST_ARCH}/cust_opapi.dll"; fi
	47	+WHEEL_PATH="${BUILD_DIR}/es_output/whl/es_custom-1.0.0-py3-none-any.whl"
	48	+require_file "${CUSTOM_OP_LIBRARY}"; require_file "${WHEEL_PATH}"
	49	+
	50	+info "Step 3/4: install the generated Python ES wheel"
	51	+python3 -m pip install --force-reinstall --upgrade --target "${BUILD_DIR}/whl_package" "${WHEEL_PATH}"
	52	+export PYTHONPATH="${BUILD_DIR}/whl_package:${SCRIPT_DIR}/src:${PYTHONPATH:-}"
	53	+export LD_LIBRARY_PATH="${BUILD_DIR}/es_output/lib64:${LD_LIBRARY_PATH:-}"
	54	+export ASCEND_CUSTOM_OPP_PATH="${BUILD_DIR}/opp:${SCRIPT_DIR}/src/ge"
	55	+export GE_PYTHON_CUSTOM_OP_SOURCE="${SCRIPT_DIR}/../cpp/add_custom_kernel/add_custom.asc"
	56	+export GE_PYTHON_CUSTOM_OP_BINARY="${KERNEL_BINARY}"
	57	+
	58	+info "Step 4/4: compare the online AnnotatedArgs and no-refresh graphs on NPU"
	59	+DEVICE_ID="${DEVICE_ID:-0}" python3 "${SCRIPT_DIR}/src/run.py" 2>&1 \| tee "${BUILD_DIR}/runtime.log"
	60	+grep -Fq "NPU_EXECUTION=PASS" "${BUILD_DIR}/runtime.log" \|\| { error "NPU execution marker is missing"; exit 1; }
	61	+for marker in "AnnotatedAddCustom precision check PASS" "NoRefreshAddCustom precision check PASS" \
	62	+ "[Perf] AnnotatedAddCustom" "[Perf] NoRefreshAddCustom" "[Perf] Annotated speedup"; do
	63	+ grep -Fq "${marker}" "${BUILD_DIR}/runtime.log" \|\| { error "Runtime marker is missing: ${marker}"; exit 1; }
	64	+done
	65	+info "Online Python custom-op pipeline PASS"

Aexamples/custom_op/annotated_args_refresh_add_custom/online/python/src/ge/annotated_add_custom.py+254-0

@@ -0,0 +1,254 @@
	1	+#!/usr/bin/env python3
	2	+# -- coding: utf-8 --
	3	+# -----------------------------------------------------------------------------------------------------------
	4	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	5	+# This program is free software; you can redistribute it and/or modify it under the terms and conditions of
	6	+# CANN Open Software License Agreement Version 2.0 (the "License").
	7	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	8	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	9	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	10	+# See LICENSE in the root of the software repository for the full text of the License.
	11	+# -----------------------------------------------------------------------------------------------------------
	12	+
	13	+"""Python prototype and AnnotatedArgs implementation for the online sample."""
	14	+
	15	+import atexit
	16	+import ctypes
	17	+import hashlib
	18	+import logging
	19	+import os
	20	+import subprocess
	21	+import tempfile
	22	+import threading
	23	+from dataclasses import dataclass
	24	+from pathlib import Path
	25	+
	26	+from ge.custom_op import (
	27	+ AnnotatedKernelLaunchInfo,
	28	+ get_compile_platform_info,
	29	+ get_declare_launch_args_ctx,
	30	+ get_execute_ctx,
	31	+ register_op,
	32	+ register_op_impl,
	33	+)
	34	+from ge.runtime import Tensor, TensorDesc
	35	+
	36	+
	37	+logging.basicConfig(level=logging.INFO, format="%(message)s")
	38	+_LOGGER = logging.getLogger(__name__)
	39	+_LOGGER.info("PY_ANNOTATED_ARGS_MODULE_LOADED=1")
	40	+
	41	+
	42	+def _infer_add(x1, x2, op_type):
	43	+ if x1.shape.origin_shape.dims != x2.shape.origin_shape.dims:
	44	+ raise ValueError("{} inputs must have the same shape".format(op_type))
	45	+ if x1.data_type != x2.data_type:
	46	+ raise ValueError("{} inputs must have the same data type".format(op_type))
	47	+ return TensorDesc(x1.shape, x1.data_type)
	48	+
	49	+
	50	+@register_op(op_type="AnnotatedAddCustom")
	51	+def annotated_add_custom_infer_meta(x1: TensorDesc, x2: TensorDesc) -> TensorDesc:
	52	+ """Infer output metadata for the declarative address-refresh operator."""
	53	+
	54	+ _LOGGER.info("PY_ANNOTATED_ARGS_INFER_META_ENTER=1")
	55	+ output0 = _infer_add(x1, x2, "AnnotatedAddCustom")
	56	+ _LOGGER.info("PY_ANNOTATED_ARGS_INFER_META_EXIT=1")
	57	+ return output0
	58	+
	59	+
	60	+@register_op(op_type="NoRefreshAddCustom")
	61	+def no_refresh_add_custom_infer_meta(x1: TensorDesc, x2: TensorDesc) -> TensorDesc:
	62	+ """Infer output metadata for the no-refresh comparison operator."""
	63	+
	64	+ _LOGGER.info("PY_NO_REFRESH_INFER_META_ENTER=1")
	65	+ output0 = _infer_add(x1, x2, "NoRefreshAddCustom")
	66	+ _LOGGER.info("PY_NO_REFRESH_INFER_META_EXIT=1")
	67	+ return output0
	68	+
	69	+
	70	+_KERNEL_NAME = "add_custom"
	71	+_BLOCK_SIZE = 1024
	72	+_KERNEL_SOURCE = Path(
	73	+ os.environ.get(
	74	+ "GE_PYTHON_CUSTOM_OP_SOURCE",
	75	+ str(
	76	+ Path(__file__).resolve().parents[3]
	77	+ / "cpp"
	78	+ / "add_custom_kernel"
	79	+ / "add_custom.asc"
	80	+ ),
	81	+ )
	82	+)
	83	+
	84	+
	85	+@dataclass(frozen=True)
	86	+class _Artifact:
	87	+ binary: bytes
	88	+ block_dim: int
	89	+
	90	+
	91	+def _dims(tensor: Tensor):
	92	+ return tuple(int(dim) for dim in tensor.storage_shape.dims)
	93	+
	94	+
	95	+def _validate(x: Tensor, y: Tensor, z: Tensor) -> None:
	96	+ if _dims(x) != _dims(y) or _dims(x) != _dims(z):
	97	+ raise ValueError("AnnotatedAddCustom requires matching tensor shapes")
	98	+ if x.data_type != y.data_type or x.data_type != z.data_type:
	99	+ raise ValueError("AnnotatedAddCustom requires matching data types")
	100	+ elements = 1
	101	+ for dim in _dims(x):
	102	+ elements *= dim
	103	+ if not _dims(x) or elements % _BLOCK_SIZE != 0:
	104	+ raise ValueError("the sample kernel requires a positive size divisible by 1024")
	105	+
	106	+
	107	+def _compile_kernel(platform) -> bytes:
	108	+ ascend_home = os.environ.get("ASCEND_HOME_PATH")
	109	+ if not ascend_home or not _KERNEL_SOURCE.is_file():
	110	+ raise RuntimeError(
	111	+ "ASCEND_HOME_PATH and the Ascend C kernel source are required"
	112	+ )
	113	+ include = Path(ascend_home) / "asc" / "include"
	114	+ out = Path(os.environ.get("GE_PYTHON_CUSTOM_OP_BUILD_DIR", tempfile.gettempdir()))
	115	+ out.mkdir(parents=True, exist_ok=True)
	116	+ digest = hashlib.sha256(_KERNEL_SOURCE.read_bytes()).hexdigest()[:16]
	117	+ host = out / ("annotated_add_custom_" + digest + ".o")
	118	+ binary = out / ("annotated_add_custom_" + digest + ".aicore.o")
	119	+ if not binary.exists():
	120	+ arch = str(platform.get_platform_resource("version", "NpuArch")).strip()
	121	+ arch = arch[4:] if arch.startswith("dav-") else arch
	122	+ subprocess.run(
	123	+ [
	124	+ "bisheng",
	125	+ "-c",
	126	+ str(_KERNEL_SOURCE),
	127	+ "-o",
	128	+ str(host),
	129	+ "--npu-arch=dav-" + arch,
	130	+ "-I" + str(include),
	131	+ ],
	132	+ check=True,
	133	+ )
	134	+ subprocess.run(
	135	+ [
	136	+ "llvm-objcopy",
	137	+ "-O",
	138	+ "binary",
	139	+ "--only-section=.aicore_binary",
	140	+ str(host),
	141	+ str(binary),
	142	+ ],
	143	+ check=True,
	144	+ )
	145	+ data = binary.read_bytes()
	146	+ if not data:
	147	+ raise RuntimeError("Ascend C kernel binary is empty")
	148	+ return bytes(data)
	149	+
	150	+
	151	+@register_op_impl(op_type="AnnotatedAddCustom")
	152	+class AnnotatedAddCustom:
	153	+ """Compile the Ascend C kernel and publish declarative launch arguments."""
	154	+
	155	+ def __init__(self):
	156	+ self._artifacts = {}
	157	+ self._lock = threading.RLock()
	158	+
	159	+ def compile(self, x: Tensor, y: Tensor, z: Tensor) -> None:
	160	+ _validate(x, y, z)
	161	+ key = (_dims(x), str(x.data_type))
	162	+ with self._lock:
	163	+ if key not in self._artifacts:
	164	+ print("PY_COMPILE_CALLBACK_ENTER=1", flush=True)
	165	+ platform = get_compile_platform_info()
	166	+ elements = 1
	167	+ for dim in _dims(x):
	168	+ elements *= dim
	169	+ self._artifacts[key] = _Artifact(
	170	+ _compile_kernel(platform), elements // _BLOCK_SIZE
	171	+ )
	172	+
	173	+ def declare_launch_args(self, x: Tensor, y: Tensor, z: Tensor) -> None:
	174	+ _validate(x, y, z)
	175	+ artifact = self._artifacts.get((_dims(x), str(x.data_type)))
	176	+ if artifact is None:
	177	+ raise RuntimeError("AnnotatedAddCustom compile cache miss")
	178	+ ctx = get_declare_launch_args_ctx()
	179	+ args = ctx.create_kernel_args()
	180	+ args.append_input(0, x)
	181	+ args.append_input(1, y)
	182	+ args.append_output(0, z)
	183	+ ctx.add_launch(
	184	+ AnnotatedKernelLaunchInfo(
	185	+ kernel_name=_KERNEL_NAME,
	186	+ kernel_bin=artifact.binary,
	187	+ block_dim=artifact.block_dim,
	188	+ stream_id=ctx.get_stream_id(),
	189	+ ),
	190	+ args,
	191	+ )
	192	+
	193	+
	194	+def _check_acl(ret, action):
	195	+ if ret != 0:
	196	+ raise RuntimeError("{} failed, ret={}".format(action, ret))
	197	+
	198	+
	199	+def _load_kernel():
	200	+ import acl
	201	+
	202	+ binary_path = Path(os.environ.get("GE_PYTHON_CUSTOM_OP_BINARY", ""))
	203	+ if not binary_path.is_file():
	204	+ raise RuntimeError("kernel binary not found: {}".format(binary_path))
	205	+ handle, ret = acl.rt.binary_load_from_file(str(binary_path), [])
	206	+ _check_acl(ret, "acl.rt.binary_load_from_file")
	207	+ try:
	208	+ function, ret = acl.rt.binary_get_function(handle, _KERNEL_NAME)
	209	+ _check_acl(ret, "acl.rt.binary_get_function")
	210	+ except Exception:
	211	+ acl.rt.binary_unload(handle)
	212	+ raise
	213	+ atexit.register(lambda: acl.rt.binary_unload(handle))
	214	+ return int(function)
	215	+
	216	+
	217	+def _launch(func_handle, x, y, z, stream, elements):
	218	+ import acl
	219	+
	220	+ args_handle, ret = acl.rt.kernel_args_init(func_handle)
	221	+ _check_acl(ret, "acl.rt.kernel_args_init")
	222	+ values = []
	223	+ for name, value in (("x", x), ("y", y), ("z", z)):
	224	+ host_value = ctypes.c_uint64(int(value))
	225	+ values.append(host_value)
	226	+ _, ret = acl.rt.kernel_args_append(
	227	+ args_handle, ctypes.addressof(host_value), ctypes.sizeof(host_value)
	228	+ )
	229	+ _check_acl(ret, "acl.rt.kernel_args_append({})".format(name))
	230	+ _check_acl(acl.rt.kernel_args_finalize(args_handle), "acl.rt.kernel_args_finalize")
	231	+ blocks = int(elements) // _BLOCK_SIZE
	232	+ _check_acl(
	233	+ acl.rt.launch_kernel_with_config(
	234	+ func_handle, blocks, stream, [], args_handle, 0
	235	+ ),
	236	+ "acl.rt.launch_kernel_with_config",
	237	+ )
	238	+ _ = values
	239	+
	240	+
	241	+@register_op_impl(op_type="NoRefreshAddCustom")
	242	+class NoRefreshAddCustom:
	duhua19 天前没有看到 NoRefreshAddCustom 的Python原型注册呢 L lfz2812 19 天前评论： 19 天前评论：
	243	+ """Execute the same Ascend C kernel through the ordinary Python path."""
	244	+
	245	+ def execute(self, x: Tensor, y: Tensor) -> None:
	246	+ if _dims(x) != _dims(y) or int(x.shape_size) % _BLOCK_SIZE != 0:
	247	+ raise ValueError(
	248	+ "NoRefreshAddCustom requires matching shapes divisible by 1024"
	249	+ )
	250	+ ctx = get_execute_ctx()
	251	+ output = ctx.malloc_output_tensor(0, x.shape, x.format, x.data_type)
	252	+ _launch(
	253	+ _load_kernel(), x.addr, y.addr, output.addr, ctx.get_stream(), x.shape_size
	254	+ )

Rexamples/custom_op/args_refresh_add_custom/python/src/run.py→examples/custom_op/annotated_args_refresh_add_custom/online/python/src/run.py+108-57

	@@ -2,7 +2,7 @@
2	# -- coding: utf-8 --	2	# -- coding: utf-8 --
3	# -----------------------------------------------------------------------------------------------------------	3	# -----------------------------------------------------------------------------------------------------------
4	# Copyright (c) 2026 Huawei Technologies Co., Ltd.	4	# Copyright (c) 2026 Huawei Technologies Co., Ltd.
5	-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of	5	+# This program is free software; you can redistribute it and/or modify it under the terms and conditions of
6	# CANN Open Software License Agreement Version 2.0 (the "License").	6	# CANN Open Software License Agreement Version 2.0 (the "License").
7	# Please refer to the License for details. You may not use this file except in compliance with the License.	7	# Please refer to the License for details. You may not use this file except in compliance with the License.
8	# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,	8	# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	@@ -10,75 +10,105 @@
10	# See LICENSE in the root of the software repository for the full text of the License.	10	# See LICENSE in the root of the software repository for the full text of the License.
11	# -----------------------------------------------------------------------------------------------------------	11	# -----------------------------------------------------------------------------------------------------------
12		12
13	-import traceback	13	+"""Compare declarative address refresh with schema-bound Execute online."""
14	-from typing import List
15		14
16	-from ge.es.graph_builder import GraphBuilder, attr_scope	15	+import os
		16	+import time
		17	+import traceback
		18	+
		19	+from ge.es.graph_builder import GraphBuilder
17	from ge.ge_global import GeApi	20	from ge.ge_global import GeApi
18	from ge.graph import Tensor	21	from ge.graph import Tensor
19	-from ge.graph.types import DataType, Format	22	+from ge.graph.types import DataType, Format, Placement
20	from ge.session import Session	23	from ge.session import Session
21		24
22	try:	25	try:
23	- from ge.es.custom import AddPythonCustomOp	26	+ from ge.es.custom import AnnotatedAddCustom, NoRefreshAddCustom
24	except ImportError as import_error:	27	except ImportError as import_error:
25	raise RuntimeError(	28	raise RuntimeError(
26	- "未找到 ge.es.custom.AddPythonCustomOp。请先执行 bash run.sh 生成并加载 es_custom Python ES API。"	29	+ "Custom-op ES APIs are unavailable. Run run.sh first."
27	) from import_error	30	) from import_error
28		31
29		32
30	-GRAPH_ID = 0	33	+ANNOTATED_GRAPH_ID = 0
31	-DEVICE_ID = 0	34	+NO_REFRESH_GRAPH_ID = 1
32	-NUM_ELEMENTS = 1024	35	+DEVICE_ID = int(os.environ.get("DEVICE_ID", "0"))
		36	+NUM_ELEMENTS = 8 * 1024
		37	+WARMUP_ITERS = 5
		38	+BENCHMARK_ITERS = 100
33		39
34		40
35	-def build_graph():	41	+def build_graph(name, op_factory):
36	- builder = GraphBuilder("add_python_graph_test")	42	+ builder = GraphBuilder(name)
37	input_x = builder.create_input(	43	input_x = builder.create_input(
38	index=0,	44	index=0,
39	name="data_x",	45	name="data_x",
40	data_type=DataType.DT_FLOAT,	46	data_type=DataType.DT_FLOAT,
41	format=Format.FORMAT_ND,	47	format=Format.FORMAT_ND,
42	- shape=[-1],	48	+ shape=[NUM_ELEMENTS],
43	)	49	)
44	input_y = builder.create_input(	50	input_y = builder.create_input(
45	index=1,	51	index=1,
46	name="data_y",	52	name="data_y",
47	data_type=DataType.DT_FLOAT,	53	data_type=DataType.DT_FLOAT,
48	format=Format.FORMAT_ND,	54	format=Format.FORMAT_ND,
49	- shape=[-1],	55	+ shape=[NUM_ELEMENTS],
50	)	56	)
51	-	57	+ output0 = op_factory(input_x, input_y)
52	- # Remove this lock after Python custom ops support infer-shape and infer-datatype registration.	58	+ builder.set_graph_output(output0, 0)
53	- # Until then, GE's fallback inference may overwrite the explicit output dtype with an undefined origin dtype.
54	- with attr_scope({"_out_shape_locked": True}):
55	- output_z = AddPythonCustomOp(input_x, input_y)
56	- output_z.set_shape([-1]).set_format(Format.FORMAT_ND).set_data_type(
57	- DataType.DT_FLOAT
58	- )
59	- builder.set_graph_output(output_z, 0)
60	return builder.build_and_reset()	59	return builder.build_and_reset()
61		60
62		61
63	-def build_input_data(start: float, step: float) -> List[float]:	62	+def build_input_data(start, step):
64	- return [start + float(i) * step for i in range(NUM_ELEMENTS)]	63	+ return [start + float(index) * step for index in range(NUM_ELEMENTS)]
65		64
66		65
67	-def build_input_tensor(data: List[float]) -> Tensor:	66	+def build_device_tensor(data):
68	- return Tensor(data, None, DataType.DT_FLOAT, Format.FORMAT_ND, [NUM_ELEMENTS])	67	+ return Tensor(
69	-	68	+ data,
70	-	69	+ None,
71	-def print_output_tensor(output: Tensor) -> None:	70	+ DataType.DT_FLOAT,
72	- print(	71	+ Format.FORMAT_ND,
73	- "[Sample] output shape={}, dtype={}, format={}".format(	72	+ [NUM_ELEMENTS],
74	- list(output.shape),	73	+ Placement.PLACEMENT_DEVICE,
75	- output.data_type,
76	- output.format,
77	- )
78	)	74	)
79		75
80		76
81	-def run_graph() -> int:	77	+def build_input_sets():
		78	+ values = (
		79	+ (build_input_data(1.0, 1.0), build_input_data(10.0, 0.5)),
		80	+ (build_input_data(3.0, 2.0), build_input_data(20.0, 0.25)),
		81	+ )
		82	+ return [([build_device_tensor(x), build_device_tensor(y)], x, y) for x, y in values]
		83	+
		84	+
		85	+def validate_graph(session, graph_id, input_sets, graph_name):
		86	+ for inputs, values_x, values_y in input_sets:
		87	+ outputs = session.run_graph(graph_id, inputs)
		88	+ if len(outputs) != 1:
		89	+ raise RuntimeError(
		90	+ "{} returned {} outputs".format(graph_name, len(outputs))
		91	+ )
		92	+ actual = outputs[0].data
		93	+ expected = [x + y for x, y in zip(values_x, values_y)]
		94	+ max_error = max(abs(value - golden) for value, golden in zip(actual, expected))
		95	+ if max_error > 1.0e-5:
		96	+ raise RuntimeError(
		97	+ "{} precision check failed, max_error={}".format(graph_name, max_error)
		98	+ )
		99	+ print("[OnlinePython] {} precision check PASS".format(graph_name))
		100	+
		101	+
		102	+def benchmark_graph(session, graph_id, input_sets):
		103	+ for iteration in range(WARMUP_ITERS):
		104	+ session.run_graph(graph_id, input_sets[iteration % len(input_sets)][0])
		105	+ start = time.perf_counter()
		106	+ for iteration in range(BENCHMARK_ITERS):
		107	+ session.run_graph(graph_id, input_sets[iteration % len(input_sets)][0])
		108	+ return (time.perf_counter() - start) * 1.0e6
		109	+
		110	+
		111	+def run_graph():
82	options = {	112	options = {
83	"ge.exec.deviceId": str(DEVICE_ID),	113	"ge.exec.deviceId": str(DEVICE_ID),
84	"ge.graphRunMode": "1",	114	"ge.graphRunMode": "1",
	@@ -86,38 +116,59 @@ def run_graph() -> int:
86	ge_api = GeApi()	116	ge_api = GeApi()
87	session = None	117	session = None
88	ge_initialized = False	118	ge_initialized = False
89	- graph_added = False	119	+ graph_ids = []
90		120
91	try:	121	try:
92	ge_api.ge_initialize(options)	122	ge_api.ge_initialize(options)
93	ge_initialized = True	123	ge_initialized = True
94	session = Session(options)	124	session = Session(options)
95	- session.add_graph(GRAPH_ID, build_graph())	125	+ session.add_graph(
96	- graph_added = True	126	+ ANNOTATED_GRAPH_ID,
97	- print("[Sample] graph added, graph_id={}".format(GRAPH_ID))	127	+ build_graph("python_annotated_graph", AnnotatedAddCustom),
		128	+ )
		129	+ graph_ids.append(ANNOTATED_GRAPH_ID)
		130	+ session.add_graph(
		131	+ NO_REFRESH_GRAPH_ID,
		132	+ build_graph("python_no_refresh_graph", NoRefreshAddCustom),
		133	+ )
		134	+ graph_ids.append(NO_REFRESH_GRAPH_ID)
		135	+ input_sets = build_input_sets()
98		136
99	- inputs = [	137	+ validate_graph(session, ANNOTATED_GRAPH_ID, input_sets, "AnnotatedAddCustom")
100	- build_input_tensor(build_input_data(1.0, 1.0)),	138	+ validate_graph(session, NO_REFRESH_GRAPH_ID, input_sets, "NoRefreshAddCustom")
101	- build_input_tensor(build_input_data(10.0, 10.0)),	139	+ annotated_us = benchmark_graph(session, ANNOTATED_GRAPH_ID, input_sets)
102	- ]	140	+ no_refresh_us = benchmark_graph(session, NO_REFRESH_GRAPH_ID, input_sets)
103	- outputs = session.run_graph(GRAPH_ID, inputs)	141	+ annotated_avg = annotated_us / BENCHMARK_ITERS
104	- if not outputs:	142	+ no_refresh_avg = no_refresh_us / BENCHMARK_ITERS
105	- raise RuntimeError("RunGraph success but outputs is empty")	143	+ speedup = no_refresh_us / annotated_us if annotated_us > 0.0 else 0.0
106		144
107	- print("[Sample] run_graph finished, outputs={}".format(len(outputs)))	145	+ print("[Perf] input shape: [{}], dtype: float32".format(NUM_ELEMENTS))
108	- print_output_tensor(outputs[0])	146	+ print("[Perf] iters: {}".format(BENCHMARK_ITERS))
		147	+ print(
		148	+ "[Perf] AnnotatedAddCustom: {:.3f} us (avg {:.3f} us/iter)".format(
		149	+ annotated_us, annotated_avg
		150	+ )
		151	+ )
		152	+ print(
		153	+ "[Perf] NoRefreshAddCustom: {:.3f} us (avg {:.3f} us/iter)".format(
		154	+ no_refresh_us, no_refresh_avg
		155	+ )
		156	+ )
		157	+ print("[Perf] Annotated speedup: {:.3f}x".format(speedup))
		158	+ print("[OnlinePython] NPU_EXECUTION=PASS")
109	return 0	159	return 0
110	except Exception as exc:	160	except Exception as exc:
111	- print("[Sample] run_graph failed: {}".format(exc))	161	+ print("[OnlinePython] run_graph failed: {}".format(exc))
112	traceback.print_exc()	162	traceback.print_exc()
113	return 1	163	return 1
114	finally:	164	finally:
115	- try:	165	+ if session is not None:
116	- if graph_added and session is not None:	166	+ for graph_id in graph_ids:
117	- session.remove_graph(GRAPH_ID)	167	+ session.remove_graph(graph_id)
118	- finally:	168	+ session = None
119	- if ge_initialized:	169	+ input_sets = []
120	- ge_api.ge_finalize()	170	+ if ge_initialized:
		171	+ ge_api.ge_finalize()
121		172
122		173
123	if __name__ == "__main__":	174	if __name__ == "__main__":

Dexamples/custom_op/annotated_args_refresh_add_custom/python/README.md+0-59

	@@ -1,59 +0,0 @@
1	-# Python 声明式地址刷新样例
2	-
3	-本样例使用 Python 构图，令 ATC 在编译期加载 Python 模块并执行 `declare_launch_args` callback，最后使用 ACL Python 在真实 NPU 上加载同一个 OM，以两套设备地址完成两轮地址刷新和数值校验。
4	-
5	-## 目录
6	-
7	-```text
8	-python
9	-├── CMakeLists.txt # 构建 custom OPP 注册库和 es_custom wheel
10	-├── run.sh # kernel、ES、AIR、ATC、ACL 两轮验证入口
11	-├── proto/ # AnnotatedAddCustom 算子原型
12	-└── src/
13	- ├── build_graph.py # Python 构图并生成 AIR
14	- ├── run_model.py # ACL Python 两轮 NPU 执行
15	- └── ge/annotated_add_custom.py # 编译期 declare_launch_args callback
16	-```
17	-
18	-## 依赖与运行
19	-
20	-需要已安装并配置的 CANN、BiSheng、LLVM `llvm-objcopy`、ATC、CMake、Python 3/pip、numpy 和可用 NPU。
21	-
22	-```bash
23	-source /path/to/cann/set_env.sh
24	-cd examples/custom_op/annotated_args_refresh_add_custom/python
25	-bash run.sh
26	-```
27	-
28	-仅可覆盖以下三个环境变量：
29	-
30	-- `ADD_CUSTOM_NPU_ARCH`：默认 `2201`，传给 BiSheng 的 `--npu-arch=dav-2201`。
31	-- `SOC_VERSION`：默认 `Ascend910B1`，传给 ATC。
32	-- `DEVICE_ID`：默认 `0`，传给 ACL Python 运行期。
33	-
34	-例如：`ADD_CUSTOM_NPU_ARCH=2201 SOC_VERSION=Ascend910B1 DEVICE_ID=0 bash run.sh`。
35	-
36	-## 产物和编译期证据
37	-
38	-脚本依次生成以下核心产物：
39	-
40	-- `build/add_custom.o`：从 Ascend C kernel 提取的 AI Core 二进制。
41	-- `build/opp/op_graph/lib/<os>/<arch>/libcust_opapi.so`：custom OPP 注册库（Windows 为 dll）。
42	-- `build/es_output/whl/es_custom-1.0.0-py3-none-any.whl`：Python ES wheel。
43	-- `build/annotated_add.air`：Python 构图生成的 AIR。
44	-- `build/annotated_add.om`：ATC 生成的 OM。
45	-
46	-另外，`build/atc.log` 保存编译日志，`build/runtime.log` 保存真实 NPU 运行日志。ATC 日志必须包含：
47	-
48	-- `PY_ANNOTATED_ARGS_MODULE_LOADED=1`：证明 Python callback 模块已经在 ATC 编译期导入。
49	-- `PY_ANNOTATED_ARGS_CALLBACK_ENTER=1`：证明 ATC 编译期进入了 `declare_launch_args` callback。
50	-
51	-Python callback 仅在 ATC 编译期执行。OM 运行期只消费已生成的 TaskDef，不会再次导入或回调 Python；因此运行日志不应出现上述两个标记。
52	-
53	-## 两轮地址刷新验证
54	-
55	-ACL Python 为 round 1 分配 `1 + 2 = 3` 的 x/y/z 设备地址，为 round 2 重新分配 `4 + 5 = 9` 的 x/y/z 设备地址。两轮中 x、y、z 的对应设备地址必须全部不同，对每轮 8192 个 `float32` 元素进行全量 `allclose` 校验。日志会输出各轮的十六进制地址、首值、期望值和最大误差，并以 `NPU_TWO_ROUND_VALIDATION=PASS` 表示通过。
56	-
57	-## callback 约束
58	-
59	-`append_input` 与 `append_output` 分别使用当前计算节点各自 input/output 实例的平铺 index；动态项展开出的实例占用连续 index。`AnnotatedArgsContext`、Tensor、workspace 和 args builder 都是 callback 期 borrowed 对象，不能逃逸到 callback 外。args builder 在 `add_launch` 后已 consumed，不能复用。

Dexamples/custom_op/annotated_args_refresh_add_custom/python/README_en.md+0-59

	@@ -1,59 +0,0 @@
1	-# Python Declarative Address Refresh Sample
2	-
3	-This sample builds a graph in Python, has ATC load the Python module and run the `declare_launch_args` callback at compile time, then uses ACL Python to load the same OM on a real NPU for two address-refresh rounds with independent device addresses and numeric validation.
4	-
5	-## Layout
6	-
7	-```text
8	-python
9	-├── CMakeLists.txt # Builds the custom OPP registry library and es_custom wheel
10	-├── run.sh # Kernel, ES, AIR, ATC, and two-round ACL validation entry point
11	-├── proto/ # AnnotatedAddCustom operator prototype
12	-└── src/
13	- ├── build_graph.py # Python graph builder that writes AIR
14	- ├── run_model.py # ACL Python two-round NPU execution
15	- └── ge/annotated_add_custom.py # Compile-time declare_launch_args callback
16	-```
17	-
18	-## Prerequisites and running
19	-
20	-Install and configure CANN, BiSheng, LLVM `llvm-objcopy`, ATC, CMake, Python 3/pip, numpy, and a usable NPU.
21	-
22	-```bash
23	-source /path/to/cann/set_env.sh
24	-cd examples/custom_op/annotated_args_refresh_add_custom/python
25	-bash run.sh
26	-```
27	-
28	-Only these three environment variables are overridable:
29	-
30	-- `ADD_CUSTOM_NPU_ARCH`: defaults to `2201` and is passed to BiSheng as `--npu-arch=dav-2201`.
31	-- `SOC_VERSION`: defaults to `Ascend910B1` and is passed to ATC.
32	-- `DEVICE_ID`: defaults to `0` and is passed to the ACL Python runtime.
33	-
34	-For example: `ADD_CUSTOM_NPU_ARCH=2201 SOC_VERSION=Ascend910B1 DEVICE_ID=0 bash run.sh`.
35	-
36	-## Outputs and compile-time evidence
37	-
38	-The script produces these core outputs:
39	-
40	-- `build/add_custom.o`: the AI Core binary extracted from the Ascend C kernel.
41	-- `build/opp/op_graph/lib/<os>/<arch>/libcust_opapi.so`: custom OPP registry library (dll on Windows).
42	-- `build/es_output/whl/es_custom-1.0.0-py3-none-any.whl`: Python ES wheel.
43	-- `build/annotated_add.air`: AIR built by the Python graph builder.
44	-- `build/annotated_add.om`: OM generated by ATC.
45	-
46	-It also writes the compile log to `build/atc.log` and the real NPU runtime log to `build/runtime.log`. The ATC log must contain:
47	-
48	-- `PY_ANNOTATED_ARGS_MODULE_LOADED=1`: the Python callback module was imported at ATC compile time.
49	-- `PY_ANNOTATED_ARGS_CALLBACK_ENTER=1`: ATC entered the `declare_launch_args` callback at compile time.
50	-
51	-The Python callback runs only at ATC compile time. At OM runtime, only the generated TaskDef is consumed; Python is neither imported nor called again, so neither marker may appear in the runtime log.
52	-
53	-## Two-round address-refresh validation
54	-
55	-ACL Python allocates x/y/z device addresses for round 1 with `1 + 2 = 3`, then newly allocates x/y/z addresses for round 2 with `4 + 5 = 9`. The corresponding x, y, and z device addresses must all differ between rounds, and all 8192 `float32` elements are checked with `allclose`. The log reports hexadecimal addresses, first value, expected value, and maximum error for each round; `NPU_TWO_ROUND_VALIDATION=PASS` signals success.
56	-
57	-## Callback constraints
58	-
59	-`append_input` and `append_output` respectively use the flattened index of their own input/output instance on the current compute node; instances expanded from a dynamic item occupy consecutive indexes. The `AnnotatedArgsContext`, Tensor, workspace, and args builder are borrowed callback-time objects and must not escape the callback. The args builder is consumed by `add_launch` and must not be reused.

Dexamples/custom_op/annotated_args_refresh_add_custom/python/src/ge/annotated_add_custom.py+0-51

	@@ -1,51 +0,0 @@
1	-#!/usr/bin/env python3
2	-# -- coding: utf-8 --
3	-# -----------------------------------------------------------------------------------------------------------
4	-# Copyright (c) 2026 Huawei Technologies Co., Ltd.
5	-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6	-# CANN Open Software License Agreement Version 2.0 (the "License").
7	-# Please refer to the License for details. You may not use this file except in compliance with the License.
8	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9	-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10	-# See LICENSE in the root of the software repository for the full text of the License.
11	-# -----------------------------------------------------------------------------------------------------------
12	-
13	-"""DLA callback for the AnnotatedAddCustom Python example."""
14	-
15	-import logging
16	-from pathlib import Path
17	-
18	-from ge.custom_op import (
19	- AnnotatedKernelLaunchInfo,
20	- get_declare_launch_args_ctx,
21	- register_op_impl,
22	-)
23	-from ge.runtime import Tensor
24	-
25	-
26	-_KERNEL_BIN_PATH = Path(__file__).resolve().parents[2] / "build" / "add_custom.o"
27	-logging.basicConfig(level=logging.INFO, format="%(message)s")
28	-_LOGGER = logging.getLogger(__name__)
29	-_LOGGER.info("PY_ANNOTATED_ARGS_MODULE_LOADED=1")
30	-_KERNEL_BIN = _KERNEL_BIN_PATH.read_bytes()
31	-
32	-
33	-@register_op_impl(op_type="AnnotatedAddCustom")
34	-class AnnotatedAddCustom:
35	- def declare_launch_args(self, x1: Tensor, x2: Tensor, y: Tensor) -> None:
36	- _ = self
37	- _LOGGER.info("PY_ANNOTATED_ARGS_CALLBACK_ENTER=1")
38	- ctx = get_declare_launch_args_ctx()
39	- args = ctx.create_kernel_args()
40	- args.append_input(0, x1)
41	- args.append_input(1, x2)
42	- args.append_output(0, y)
43	- ctx.add_launch(
44	- AnnotatedKernelLaunchInfo(
45	- kernel_name="add_custom",
46	- kernel_bin=_KERNEL_BIN,
47	- block_dim=8,
48	- stream_id=ctx.get_stream_id(),
49	- ),
50	- args,
51	- )

Dexamples/custom_op/args_refresh_add_custom/python/CMakeLists.txt+0-51

	@@ -1,51 +0,0 @@
1	-cmake_minimum_required(VERSION 3.16)
2	-project(args_refresh_add_python_custom LANGUAGES CXX)
3	-
4	-set(CMAKE_CXX_STANDARD 17)
5	-set(CMAKE_CXX_STANDARD_REQUIRED ON)
6	-set(CMAKE_CXX_EXTENSIONS OFF)
7	-
8	-if(NOT CMAKE_BUILD_TYPE)
9	- set(CMAKE_BUILD_TYPE Release CACHE STRING "Build type" FORCE)
10	-endif()
11	-
12	-set(ES_OUTPUT_PATH "${CMAKE_BINARY_DIR}/es_output")
13	-
14	-set(ASCEND_HOME_PATH_OVERRIDE "" CACHE PATH "Optional ASCEND_HOME_PATH override")
15	-if(ASCEND_HOME_PATH_OVERRIDE)
16	- set(ASCEND_HOME_PATH "${ASCEND_HOME_PATH_OVERRIDE}")
17	-else()
18	- set(ASCEND_HOME_PATH "$ENV{ASCEND_HOME_PATH}")
19	-endif()
20	-
21	-if(ASCEND_HOME_PATH)
22	- message(STATUS "ASCEND_HOME_PATH: ${ASCEND_HOME_PATH}")
23	-else()
24	- message(FATAL_ERROR "ASCEND_HOME_PATH is empty. Please source CANN set_env.sh first.")
25	-endif()
26	-
27	-list(APPEND CMAKE_MODULE_PATH "${ASCEND_HOME_PATH}/include/ge/cmake")
28	-find_package(GenerateEsPackage REQUIRED)
29	-
30	-add_compile_definitions(
31	- _GLIBCXX_USE_CXX11_ABI=0
32	- google=ascend_private
33	-)
34	-
35	-add_library(add_python_custom_op_proto SHARED
36	- proto/add_custom.cc
37	-)
38	-target_compile_options(add_python_custom_op_proto PRIVATE
39	- -fvisibility=hidden
40	-)
41	-target_compile_definitions(add_python_custom_op_proto PRIVATE OP_PROTO_LIB)
42	-target_include_directories(add_python_custom_op_proto PRIVATE
43	- "${ASCEND_HOME_PATH}/include"
44	- "${ASCEND_HOME_PATH}/include/external"
45	-)
46	-
47	-add_es_library_and_whl(
48	- ES_LINKABLE_AND_ALL_TARGET es_custom
49	- OPP_PROTO_TARGET add_python_custom_op_proto
50	- OUTPUT_PATH ${ES_OUTPUT_PATH}
51	-)

Dexamples/custom_op/args_refresh_add_custom/python/README.md+0-94

	@@ -1,94 +0,0 @@
1	-# ArgsUpdater Add Custom Python 实现样例
2	-
3	-本目录是 `examples/custom_op/args_refresh_add_custom/python` 的最小 Python 自定义算子执行样例，用于验证 GE 可以加载 Python 自定义算子，并通过 Python ES API 构图和执行。
4	-
5	-## 范围
6	-
7	-- 自定义算子原型仍使用 C++ `REG_OP` 注册：`AddPythonCustomOp`
8	-- 算子实现使用 schema-bound Python `execute(x, y)`，`x`、`y` 按 `REG_OP` canonical IR 顺序绑定
9	-- 通过 `gen_esb` 生成 `ge.es.custom.AddPythonCustomOp` Python ES API，并使用 Python `GraphBuilder` 和 `Session.run_graph` 构图执行
10	-- Ascend C kernel 复用 `../cpp/add_custom_kernel/add_custom.asc`
11	-- `run.sh` 通过 `bisheng` 将 Ascend C kernel 预编译为 host object，并从 `.aicore_binary` section 提取 AI Core device binary
12	-- `execute` 中通过 `get_execute_ctx()` 获取仅在回调期间有效的执行上下文，用于分配输出和获取 stream；随后通过 ACL Python runtime 加载提取出的 device binary、使用 `kernel_args_*` 准备 x/y/z 三个地址参数，并调用 `acl.rt.launch_kernel_with_config`
13	-- 不做 ArgsUpdater 地址刷新优化、性能对比或精度校验
14	-
15	-## 目录
16	-
17	-```text
18	-args_refresh_add_custom
19	-├── cpp
20	-│ ├── add_custom_kernel
21	-│ │ └── add_custom.asc # 本样例复用的 Ascend C kernel
22	-└── python
23	- ├── CMakeLists.txt # 通过 gen_esb 生成 es_custom wheel
24	- ├── run.sh
25	- ├── proto
26	- │ ├── add_custom.h # REG_OP 原型定义，供 gen_esb 生成 Python ES API
27	- │ └── add_custom.cc # 自定义算子 proto 编译入口
28	- └── src
29	- ├── run.py # Python 构图和执行入口
30	- └── ge
31	- └── add_custom.py # Python 自定义算子实现
32	-```
33	-
34	-## 前置条件
35	-
36	-- 参考 [安装指导](../../../../docs/zh/quick_install.md) 正确安装`toolkit`和`ops`包
37	-- 设置环境变量（假设包安装在 /usr/local/Ascend/）
38	-```
39	-source /usr/local/Ascend/cann/set_env.sh
40	-```
41	-- run 包编译使用的 Python 版本与执行本样例的 Python 版本一致。当前 Python 自定义算子加载链路还不支持跨 Python 版本兼容
42	-- 当前 Python 环境可导入 `ge.custom_op` 和 `acl`
43	-- CANN run 包支持 schema-bound Python 自定义算子 `execute(inputs, *attrs)` 和 `get_execute_ctx()` 接口
44	-
45	-## Conda 环境示例（Python 3.11）
46	-
47	-如果本机没有现成的匹配环境，可以参考下面的方式创建：
48	-
49	-```bash
50	-conda create -n ge-custom-op-py311 python=3.11 -y
51	-conda activate ge-custom-op-py311
52	-python -m pip install --upgrade pip
53	-python -m pip install attrs decorator sympy numpy psutil scipy
54	-```
55	-
56	-创建环境后，请确认：
57	-
58	-- 该环境中的 Python 版本与 run 包编译时使用的 Python 版本一致
59	-- 再执行 `source ${ASCEND_HOME_PATH}/set_env.sh` 完成 CANN 环境变量设置
60	-- 最后按本文“运行”章节执行样例
61	-
62	-## 运行
63	-
64	-`run.sh` 会先通过 `bisheng` 编译 `build/add_custom.host.o`，再通过 `llvm-objcopy` 提取 `build/add_custom.aicore.o`，然后基于 `proto/add_custom.h` 生成 `es_custom` wheel，使当前 Python 进程可以导入 `ge.es.custom.AddPythonCustomOp`。
65	-`ADD_CUSTOM_NPU_ARCH` 对应 Bisheng `--npu-arch=dav-xxxx` 中的 `xxxx`，默认 `2201`。可在官方文档中查询 [AI 处理器型号和 `__NPU_ARCH__` 的对应关系](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/910beta3/compiler/BishengCompiler/atlas_bisheng_10_0005.html#ZH-CN_TOPIC_0000002594810544__table547713114286)后指定运行，例如 `ADD_CUSTOM_NPU_ARCH=2202 bash run.sh`。
66	-
67	-```bash
68	-source ${ASCEND_HOME_PATH}/set_env.sh
69	-cd examples/custom_op/args_refresh_add_custom/python
70	-bash run.sh
71	-```
72	-
73	-成功时可以看到类似输出：
74	-
75	-```text
76	-[Sample] graph added, graph_id=0
77	-[PythonCustomOp] loaded kernel binary=.../build/add_custom.aicore.o, kernel=add_custom
78	-[PythonCustomOp] AddPythonCustomOp.execute(x, y) called
79	-[PythonCustomOp] x shape=[1024], dtype=0, addr=0x...
80	-[PythonCustomOp] y shape=[1024], dtype=0, addr=0x...
81	-[PythonCustomOp] z shape=[1024], dtype=0, addr=0x...
82	-[PythonCustomOp] stream=0x...
83	-[PythonCustomOp] kernel args handle=0x..., num_blocks=1
84	-[PythonCustomOp] acl.rt.launch_kernel_with_config ret=0
85	-[Sample] run_graph finished, outputs=1
86	-[Sample] output shape=[1024], dtype=0, format=2
87	-```
88	-
89	-## 说明
90	-
91	-`src/run.py` 使用 `GraphBuilder` 创建两个 `Data` 输入，通过生成的 `ge.es.custom.AddPythonCustomOp` 构图，显式设置输出 shape/data type/format 后调用 `Session.run_graph` 执行。
92	-`run.sh` 先通过 Bisheng 将 Ascend C 源码编译为 `add_custom.host.o`，再通过 `llvm-objcopy --only-section=.aicore_binary` 提取 `add_custom.aicore.o`。`AddPythonCustomOp` 使用 Python 完成 host 侧调度：读取输入/输出地址、通过 `acl.rt.binary_load_from_file` 加载提取后的 `add_custom.aicore.o`。当前 `binary_load_from_file` 不支持 `ACL_RT_LOAD_BINARY_OPT_MAGIC`，因此加载选项传空列表。
93	-kernel 参数通过 `acl.rt.kernel_args_init`、`acl.rt.kernel_args_append`、`acl.rt.kernel_args_finalize` 按 x/y/z 顺序追加，并通过 ACL Python runtime 的 `acl.rt.launch_kernel_with_config` 下发 `add_custom` kernel。`launch_kernel_with_config` 的 `cfg` 传空列表，使用 runtime 默认配置。
94	-当前阶段 C++ 代码只承担 `REG_OP` 原型声明和复用的 Ascend C kernel 源码；真正的算子执行入口在 `src/ge/add_custom.py` 的 `execute(x, y)` 中。`x`、`y` 由 GE 根据 canonical IR 自动绑定，输出分配和 stream 获取通过回调期间的 `get_execute_ctx()` 完成。

Dexamples/custom_op/args_refresh_add_custom/python/README_en.md+0-94

	@@ -1,94 +0,0 @@
1	-# ArgsUpdater Add Custom Python Implementation Sample
2	-
3	-This directory is the minimal Python custom operator execution sample at `examples/custom_op/args_refresh_add_custom/python`. It verifies that GE can load a Python custom operator and build and execute a graph through the Python ES API.
4	-
5	-## Scope
6	-
7	-- The custom operator prototype still uses the C++ `REG_OP` registration: `AddPythonCustomOp`
8	-- The operator implementation uses schema-bound Python `execute(x, y)`, with `x` and `y` bound in canonical `REG_OP` IR order
9	-- The `ge.es.custom.AddPythonCustomOp` Python ES API is generated through `gen_esb`, and the Python `GraphBuilder` and `Session.run_graph` are used for graph construction and execution
10	-- The Ascend C kernel reuses `../cpp/add_custom_kernel/add_custom.asc`
11	-- `run.sh` precompiles the Ascend C kernel into a host object through `bisheng` and extracts the AI Core device binary from the `.aicore_binary` section
12	-- In `execute`, `get_execute_ctx()` obtains the callback execution context, which is valid only during the callback, to allocate the output and obtain the stream. The extracted device binary is then loaded through the ACL Python runtime, the x/y/z address parameters are prepared using `kernel_args_*`, and `acl.rt.launch_kernel_with_config` is called
13	-- The ArgsUpdater address refresh optimization, performance comparison, and accuracy verification are not performed
14	-
15	-## Directory
16	-
17	-```text
18	-args_refresh_add_custom
19	-├── cpp
20	-│ ├── add_custom_kernel
21	-│ │ └── add_custom.asc # Ascend C kernel reused in this sample
22	-└── python
23	- ├── CMakeLists.txt # Generates the es_custom wheel through gen_esb
24	- ├── run.sh
25	- ├── proto
26	- │ ├── add_custom.h # REG_OP prototype definition for gen_esb to generate the Python ES API
27	- │ └── add_custom.cc # Custom operator proto compilation entry
28	- └── src
29	- ├── run.py # Python graph construction and execution entry
30	- └── ge
31	- └── add_custom.py # Python custom operator implementation
32	-```
33	-
34	-## Prerequisites
35	-
36	-- Follow the [Installation Guide](../../../../docs/en/quick_install.md) to install the `toolkit` and `ops` packages.
37	-- Configure the environment variables. The following example assumes that the packages are installed in `/usr/local/Ascend/`:
38	-```
39	-source /usr/local/Ascend/cann/set_env.sh
40	-```
41	-- The Python version used to compile the run package matches the Python version used to run this sample. The current Python custom operator loading path does not support cross-Python-version compatibility
42	-- The current Python environment can import `ge.custom_op` and `acl`
43	-- The CANN run package supports schema-bound Python custom operator `execute(inputs, *attrs)` and the `get_execute_ctx()` API
44	-
45	-## Conda Environment Sample (Python 3.11)
46	-
47	-If you do not have a matching environment on the local machine, you can create one as follows:
48	-
49	-```bash
50	-conda create -n ge-custom-op-py311 python=3.11 -y
51	-conda activate ge-custom-op-py311
52	-python -m pip install --upgrade pip
53	-python -m pip install attrs decorator sympy numpy psutil scipy
54	-```
55	-
56	-After creating the environment, confirm the following:
57	-
58	-- The Python version in this environment matches the Python version used to compile the run package
59	-- Run `source ${ASCEND_HOME_PATH}/set_env.sh` to configure the CANN environment variables
60	-- Follow the "Running" section in this document to run the sample
61	-
62	-## Running
63	-
64	-`run.sh` first compiles `build/add_custom.host.o` through `bisheng`, then extracts `build/add_custom.aicore.o` through `llvm-objcopy`, and then generates the `es_custom` wheel based on `proto/add_custom.h` so that the current Python process can import `ge.es.custom.AddPythonCustomOp`.
65	-`ADD_CUSTOM_NPU_ARCH` corresponds to `xxxx` in the Bisheng `--npu-arch=dav-xxxx` option, and the default value is `2201`. You can query the [mapping between AI processor models and `__NPU_ARCH__`](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/910beta3/compiler/BishengCompiler/atlas_bisheng_10_0005.html#ZH-CN_TOPIC_0000002594810544__table547713114286) in the official documentation and specify the value at runtime, for example, `ADD_CUSTOM_NPU_ARCH=2202 bash run.sh`.
66	-
67	-```bash
68	-source ${ASCEND_HOME_PATH}/set_env.sh
69	-cd examples/custom_op/args_refresh_add_custom/python
70	-bash run.sh
71	-```
72	-
73	-Upon success, you can see output similar to the following:
74	-
75	-```text
76	-[Sample] graph added, graph_id=0
77	-[PythonCustomOp] loaded kernel binary=.../build/add_custom.aicore.o, kernel=add_custom
78	-[PythonCustomOp] AddPythonCustomOp.execute(x, y) called
79	-[PythonCustomOp] x shape=[1024], dtype=0, addr=0x...
80	-[PythonCustomOp] y shape=[1024], dtype=0, addr=0x...
81	-[PythonCustomOp] z shape=[1024], dtype=0, addr=0x...
82	-[PythonCustomOp] stream=0x...
83	-[PythonCustomOp] kernel args handle=0x..., num_blocks=1
84	-[PythonCustomOp] acl.rt.launch_kernel_with_config ret=0
85	-[Sample] run_graph finished, outputs=1
86	-[Sample] output shape=[1024], dtype=0, format=2
87	-```
88	-
89	-## Description
90	-
91	-`src/run.py` uses `GraphBuilder` to create two `Data` inputs, builds the graph through the generated `ge.es.custom.AddPythonCustomOp`, explicitly sets the output shape, data type, and format, and then calls `Session.run_graph` for execution.
92	-`run.sh` first compiles the Ascend C source code into `add_custom.host.o` through Bisheng, and then extracts `add_custom.aicore.o` through `llvm-objcopy --only-section=.aicore_binary`. `AddPythonCustomOp` uses Python to perform host-side scheduling: it reads the input and output addresses and loads the extracted `add_custom.aicore.o` through `acl.rt.binary_load_from_file`. Currently, `binary_load_from_file` does not support `ACL_RT_LOAD_BINARY_OPT_MAGIC`, so an empty list is passed for the loading options.
93	-The kernel parameters are appended in the x/y/z order through `acl.rt.kernel_args_init`, `acl.rt.kernel_args_append`, and `acl.rt.kernel_args_finalize`, and the `add_custom` kernel is launched through `acl.rt.launch_kernel_with_config` of the ACL Python runtime. An empty list is passed as the `cfg` parameter of `launch_kernel_with_config` to use the default runtime configuration.
94	-At the current stage, the C++ code only provides the `REG_OP` prototype declaration and the reused Ascend C kernel source code. The actual operator execution entry is `execute(x, y)` in `src/ge/add_custom.py`. GE binds `x` and `y` from the canonical IR, while output allocation and stream access use `get_execute_ctx()` during the callback.

Dexamples/custom_op/args_refresh_add_custom/python/run.sh+0-118

	@@ -1,118 +0,0 @@
1	-#!/usr/bin/env bash
2	-# -----------------------------------------------------------------------------------------------------------
3	-# Copyright (c) 2026 Huawei Technologies Co., Ltd.
4	-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
5	-# 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,
8	-# INCLUDING 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	-set -euo pipefail
13	-
14	-SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
15	-BUILD_DIR="${SCRIPT_DIR}/build"
16	-PY_CUSTOM_OP_DIR="${SCRIPT_DIR}/src/ge"
17	-ES_WHL_PATH="${BUILD_DIR}/es_output/whl/es_custom-1.0.0-py3-none-any.whl"
18	-ES_LIB_DIR="${BUILD_DIR}/es_output/lib64"
19	-ES_WHL_INSTALL_DIR="${BUILD_DIR}/whl_package"
20	-HOST_OS="$(uname -s \| tr '[:upper:]' '[:lower:]')"
21	-HOST_ARCH="$(uname -m)"
22	-case "${HOST_ARCH}" in
23	- arm64)
24	- HOST_ARCH="aarch64"
25	- ;;
26	- amd64)
27	- HOST_ARCH="x86_64"
28	- ;;
29	-esac
30	-CUSTOM_OPP_ROOT="${BUILD_DIR}/custom_op_package"
31	-CUSTOM_OP_LIB_DIR="${CUSTOM_OPP_ROOT}/op_graph/lib/${HOST_OS}/${HOST_ARCH}"
32	-CUSTOM_OP_PROTO_SOURCE="${BUILD_DIR}/libadd_python_custom_op_proto.so"
33	-CUSTOM_OP_PROTO_TARGET="${CUSTOM_OP_LIB_DIR}/libadd_python_custom_op_proto.so"
34	-KERNEL_SOURCE_PATH="${SCRIPT_DIR}/../cpp/add_custom_kernel/add_custom.asc"
35	-KERNEL_HOST_OBJECT_PATH="${BUILD_DIR}/add_custom.host.o"
36	-KERNEL_BINARY_PATH="${BUILD_DIR}/add_custom.aicore.o"
37	-ADD_CUSTOM_NPU_ARCH="${ADD_CUSTOM_NPU_ARCH:-2201}"
38	-
39	-info() {
40	- echo "[INFO] $*"
41	-}
42	-
43	-error() {
44	- echo "[ERROR] $*" >&2
45	-}
46	-
47	-detect_jobs() {
48	- if command -v nproc >/dev/null 2>&1; then
49	- nproc
50	- return
51	- fi
52	- echo 8
53	-}
54	-
55	-if [[ -z "${ASCEND_HOME_PATH:-}" ]]; then
56	- error "ASCEND_HOME_PATH is empty. Please source CANN set_env.sh first."
57	- exit 1
58	-fi
59	-if ! command -v bisheng >/dev/null 2>&1; then
60	- error "bisheng was not found. Please source CANN set_env.sh first."
61	- exit 1
62	-fi
63	-if ! command -v llvm-objcopy >/dev/null 2>&1; then
64	- error "llvm-objcopy was not found. Please source CANN set_env.sh first."
65	- exit 1
66	-fi
67	-
68	-if [[ ! -f "${KERNEL_SOURCE_PATH}" ]]; then
69	- error "Kernel source was not found: ${KERNEL_SOURCE_PATH}"
70	- exit 1
71	-fi
72	-
73	-mkdir -p "${BUILD_DIR}"
74	-
75	-info "Step 1/4: compile Ascend C kernel and extract AI Core binary"
76	-bisheng -c "${KERNEL_SOURCE_PATH}" -o "${KERNEL_HOST_OBJECT_PATH}" --npu-arch="dav-${ADD_CUSTOM_NPU_ARCH}"
77	-llvm-objcopy -O binary --only-section=.aicore_binary "${KERNEL_HOST_OBJECT_PATH}" "${KERNEL_BINARY_PATH}"
78	-
79	-if [[ ! -s "${KERNEL_BINARY_PATH}" ]]; then
80	- error "AI Core kernel binary was not generated: ${KERNEL_BINARY_PATH}"
81	- exit 1
82	-fi
83	-info "Kernel host object=${KERNEL_HOST_OBJECT_PATH}"
84	-info "Kernel AI Core binary=${KERNEL_BINARY_PATH}"
85	-
86	-info "Step 2/4: configure and build Python ES API"
87	-cmake -S "${SCRIPT_DIR}" -B "${BUILD_DIR}" -DCMAKE_BUILD_TYPE=Release
88	-cmake --build "${BUILD_DIR}" --target build_es_custom -j"$(detect_jobs)"
89	-
90	-if [[ ! -f "${ES_WHL_PATH}" ]]; then
91	- error "es_custom wheel was not generated: ${ES_WHL_PATH}"
92	- exit 1
93	-fi
94	-if [[ ! -f "${CUSTOM_OP_PROTO_SOURCE}" ]]; then
95	- error "Custom op proto library was not generated: ${CUSTOM_OP_PROTO_SOURCE}"
96	- exit 1
97	-fi
98	-
99	-# 当前需要通过 op_graph 提前加载并持有 C++ 原型 SO。后续本样例改为由 Python 注册算子原型后，
100	-# 可以删除这段原型 SO 复制逻辑以及 ASCEND_CUSTOM_OPP_PATH 中的 CUSTOM_OPP_ROOT。
101	-cmake -E make_directory "${CUSTOM_OP_LIB_DIR}"
102	-cmake -E copy_if_different "${CUSTOM_OP_PROTO_SOURCE}" "${CUSTOM_OP_PROTO_TARGET}"
103	-info "Custom op proto library=${CUSTOM_OP_PROTO_TARGET}"
104	-
105	-info "Step 3/4: install generated es_custom Python package"
106	-python3 -m pip install --force-reinstall --upgrade --target "${ES_WHL_INSTALL_DIR}" "${ES_WHL_PATH}"
107	-
108	-export ASCEND_CUSTOM_OPP_PATH="${CUSTOM_OPP_ROOT}:${PY_CUSTOM_OP_DIR}${ASCEND_CUSTOM_OPP_PATH:+:${ASCEND_CUSTOM_OPP_PATH}}"
109	-export PYTHONPATH="${ES_WHL_INSTALL_DIR}${PYTHONPATH:+:${PYTHONPATH}}"
110	-export LD_LIBRARY_PATH="${ES_LIB_DIR}${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}"
111	-info "ASCEND_CUSTOM_OPP_PATH=${ASCEND_CUSTOM_OPP_PATH}"
112	-info "PYTHONPATH=${PYTHONPATH}"
113	-info "LD_LIBRARY_PATH=${LD_LIBRARY_PATH}"
114	-
115	-info "Step 4/4: run Python Session::run_graph sample"
116	-python3 "${SCRIPT_DIR}/src/run.py"
117	-
118	-info "Python session sample finished."

Dexamples/custom_op/args_refresh_add_custom/python/src/ge/add_custom.py+0-161

	@@ -1,161 +0,0 @@
1	-#!/usr/bin/env python3
2	-# -- coding: utf-8 --
3	-# -----------------------------------------------------------------------------------------------------------
4	-# Copyright (c) 2026 Huawei Technologies Co., Ltd.
5	-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6	-# CANN Open Software License Agreement Version 2.0 (the "License").
7	-# Please refer to the License for details. You may not use this file except in compliance with the License.
8	-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
9	-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
10	-# See LICENSE in the root of the software repository for the full text of the License.
11	-# -----------------------------------------------------------------------------------------------------------
12	-
13	-import atexit
14	-import ctypes
15	-from pathlib import Path
16	-
17	-from ge.custom_op import get_execute_ctx, register_op_impl
18	-
19	-
20	-_KERNEL_NAME = "add_custom"
21	-_KERNEL_BLOCK_SIZE = 1024
22	-
23	-
24	-def _format_addr(addr: int) -> str:
25	- return "0x{:x}".format(addr)
26	-
27	-
28	-def _shape_dims(tensor):
29	- return tensor.storage_shape.dims
30	-
31	-
32	-def _check_ret(ret, action: str) -> None:
33	- if ret != 0:
34	- raise RuntimeError("{} failed, ret={}".format(action, ret))
35	-
36	-
37	-def _get_kernel_binary_path() -> Path:
38	- kernel_binary_path = (
39	- Path(__file__).resolve().parents[2] / "build" / "add_custom.aicore.o"
40	- )
41	- if not kernel_binary_path.is_file():
42	- raise RuntimeError("kernel binary not found: {}".format(kernel_binary_path))
43	- return kernel_binary_path
44	-
45	-
46	-def _unload_binary(bin_handle: int) -> None:
47	- import acl
48	-
49	- print("[PythonCustomOp] unload kernel binary")
50	- ret = acl.rt.binary_unload(bin_handle)
51	- _check_ret(ret, "acl.rt.binary_unload")
52	-
53	-
54	-def _load_kernel():
55	- import acl
56	-
57	- kernel_binary_path = _get_kernel_binary_path()
58	- bin_handle, ret = acl.rt.binary_load_from_file(str(kernel_binary_path), [])
59	- _check_ret(ret, "acl.rt.binary_load_from_file")
60	- try:
61	- func_handle, ret = acl.rt.binary_get_function(bin_handle, _KERNEL_NAME)
62	- _check_ret(ret, "acl.rt.binary_get_function")
63	- except Exception:
64	- acl.rt.binary_unload(bin_handle)
65	- raise
66	-
67	- atexit.register(_unload_binary, bin_handle)
68	- print(
69	- "[PythonCustomOp] loaded kernel binary={}, kernel={}".format(
70	- kernel_binary_path, _KERNEL_NAME
71	- )
72	- )
73	- return int(func_handle)
74	-
75	-
76	-def _append_kernel_arg(acl, args_handle: int, value: int, name: str):
77	- host_value = ctypes.c_uint64(int(value))
78	- param_handle, ret = acl.rt.kernel_args_append(
79	- args_handle, ctypes.addressof(host_value), ctypes.sizeof(host_value)
80	- )
81	- _check_ret(ret, "acl.rt.kernel_args_append({})".format(name))
82	- _ = param_handle
83	- return host_value
84	-
85	-
86	-def _build_kernel_args(func_handle: int, x_addr: int, y_addr: int, z_addr: int):
87	- import acl
88	-
89	- args_handle, ret = acl.rt.kernel_args_init(func_handle)
90	- _check_ret(ret, "acl.rt.kernel_args_init")
91	- host_values = []
92	- for name, value in (("x", x_addr), ("y", y_addr), ("z", z_addr)):
93	- host_values.append(_append_kernel_arg(acl, args_handle, value, name))
94	- ret = acl.rt.kernel_args_finalize(args_handle)
95	- _check_ret(ret, "acl.rt.kernel_args_finalize")
96	- return args_handle, host_values
97	-
98	-
99	-def _get_num_blocks(input_x) -> int:
100	- element_count = int(input_x.shape_size)
101	- if element_count % _KERNEL_BLOCK_SIZE != 0:
102	- raise RuntimeError(
103	- "reused add_custom kernel requires element count to be a multiple of {}, got {}".format(
104	- _KERNEL_BLOCK_SIZE, element_count
105	- )
106	- )
107	- return element_count // _KERNEL_BLOCK_SIZE
108	-
109	-
110	-def _print_tensor_info(name: str, tensor) -> None:
111	- print(
112	- "[PythonCustomOp] {} shape={}, dtype={}, addr={}".format(
113	- name, _shape_dims(tensor), tensor.data_type, _format_addr(tensor.addr)
114	- )
115	- )
116	-
117	-
118	-def _launch_kernel(
119	- func_handle: int, num_blocks: int, stream: int, args_handle: int
120	-) -> None:
121	- import acl
122	-
123	- ret = acl.rt.launch_kernel_with_config(
124	- func_handle,
125	- num_blocks,
126	- stream,
127	- [],
128	- args_handle,
129	- 0,
130	- )
131	- print("[PythonCustomOp] acl.rt.launch_kernel_with_config ret={}".format(ret))
132	- if ret != 0:
133	- raise RuntimeError(
134	- "acl.rt.launch_kernel_with_config failed, ret={}".format(ret)
135	- )
136	-
137	-
138	-@register_op_impl(op_type="AddPythonCustomOp")
139	-class AddPythonCustomOp:
140	- def execute(self, x, y) -> None:
141	- ctx = get_execute_ctx()
142	- output_z = ctx.malloc_output_tensor(0, x.shape, x.format, x.data_type)
143	- num_blocks = _get_num_blocks(x)
144	- func_handle = _load_kernel()
145	-
146	- print("[PythonCustomOp] AddPythonCustomOp.execute(x, y) called")
147	- _print_tensor_info("x", x)
148	- _print_tensor_info("y", y)
149	- _print_tensor_info("z", output_z)
150	- stream = ctx.get_stream()
151	- print("[PythonCustomOp] stream={}".format(_format_addr(stream)))
152	- args_handle, host_values = _build_kernel_args(
153	- func_handle, int(x.addr), int(y.addr), int(output_z.addr)
154	- )
155	- print(
156	- "[PythonCustomOp] kernel args handle={}, num_blocks={}".format(
157	- _format_addr(args_handle), num_blocks
158	- )
159	- )
160	- _launch_kernel(func_handle, num_blocks, stream, args_handle)
161	- _ = host_values

Mruntime/v1/executor/ge_executor.cc+1-0

	@@ -340,6 +340,7 @@ Status GeExecutor::FinalizeEx() {
340	ProfilingProperties::Instance().ClearProperties();	340	ProfilingProperties::Instance().ClearProperties();
341	}	341	}
342		342
		343	+ (void)custom_op::UnloadCustomOps();
343	CustomOpSoLoader::Finalize();	344	CustomOpSoLoader::Finalize();
344	OpsKernelExecutorManager::GetInstance().Finalize();	345	OpsKernelExecutorManager::GetInstance().Finalize();
345	HostMemManager::Instance().Finalize();	346	HostMemManager::Instance().Finalize();