feat(custom_op): 新增 TileLang 自定义算子入图/在线编译/离线模型下沉样例-ge-AtomGit

Mexamples/custom_op/README.md+3-0

	@@ -13,6 +13,9 @@
13	\| `args_refresh_add_custom` \| ArgsUpdater 地址刷新 + MallocReadOnlyDevArgs + 性能对比 \| GE 在线执行 \| Ascend C \| RTC 运行时编译 \| 在线地址刷新性能对比 \| [README](./args_refresh_add_custom/cpp/README.md) \|	13	\| `args_refresh_add_custom` \| ArgsUpdater 地址刷新 + MallocReadOnlyDevArgs + 性能对比 \| GE 在线执行 \| Ascend C \| RTC 运行时编译 \| 在线地址刷新性能对比 \| [README](./args_refresh_add_custom/cpp/README.md) \|
14	\| `annotated_args_refresh_add_custom` \| AnnotatedArgsOp 声明式地址刷新在线场景性能对比+离线场景 \| GE 在线执行 + ATC 离线编译 \| Ascend C \| RTC 运行时编译 \| 支持在线性能对比和 OM 模型下沉 \| [README](./annotated_args_refresh_add_custom/README.md) \|	14	\| `annotated_args_refresh_add_custom` \| AnnotatedArgsOp 声明式地址刷新在线场景性能对比+离线场景 \| GE 在线执行 + ATC 离线编译 \| Ascend C \| RTC 运行时编译 \| 支持在线性能对比和 OM 模型下沉 \| [README](./annotated_args_refresh_add_custom/README.md) \|
15	\| `args_refresh_add_custom（Python 版本）` \| Python EagerExecuteOp 执行 \| GE 在线执行 \| Ascend C \| Bisheng 预编译 \| 不涉及 \| [README](./args_refresh_add_custom/python/README.md) \|	15	\| `args_refresh_add_custom（Python 版本）` \| Python EagerExecuteOp 执行 \| GE 在线执行 \| Ascend C \| Bisheng 预编译 \| 不涉及 \| [README](./args_refresh_add_custom/python/README.md) \|
		16	+\| `tilelang_add_custom` \| TileLang 算子通过 GE 入图 \| GE 原生 (Session API) \| TileLang \| TileLang 预编译产出 `.so` \| 不涉及 \| [README](./tilelang_add_custom/README.md) \|
		17	+\| `tilelang_add_custom_online` \| TileLang 算子在线编译 + 在线执行 \| GE 原生 (Session API) \| TileLang \| GE 编译阶段 `CompilableOp::Compile` subprocess 调用 Python 编译器 \| 不涉及 \| [README](./tilelang_add_custom_online/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) \|
16		19
17	## 通用开发流程	20	## 通用开发流程
18		21

Mexamples/custom_op/README_en.md+3-0

	@@ -13,6 +13,9 @@ This directory provides samples related to custom operator graph integration, co
13	\| `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) \|	13	\| `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	\| `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) \|	14	\| `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	\| `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) \|	15	\| `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) \|
		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) \|
		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) \|
		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) \|
16		19
17	## General Development Process	20	## General Development Process
18		21

Aexamples/custom_op/tilelang_add_custom/CMakeLists.txt+149-0

@@ -0,0 +1,149 @@
	1	+cmake_minimum_required(VERSION 3.14)
	2	+project(tilelang_add_custom LANGUAGES CXX)
	3	+
	4	+option(TILELANG_BUILD_CUSTOM_OP "Build libcust_opapi.so" ON)
	5	+option(TILELANG_BUILD_SESSION_RUN "Build session_run" ON)
	6	+
	7	+set(CMAKE_CXX_STANDARD 17)
	8	+set(CMAKE_CXX_STANDARD_REQUIRED ON)
	9	+set(CMAKE_CXX_EXTENSIONS OFF)
	10	+
	11	+if(NOT CMAKE_BUILD_TYPE)
	12	+ set(CMAKE_BUILD_TYPE Release CACHE STRING "Build type" FORCE)
	13	+endif()
	14	+
	15	+set(COMMON_COMPILE_OPTIONS
	16	+ -Wall
	17	+ -Wextra
	18	+ -Wno-unused-parameter
	19	+)
	20	+
	21	+set(PROJECT_OUTPUT_DIR "${CMAKE_SOURCE_DIR}/output")
	22	+file(MAKE_DIRECTORY "${PROJECT_OUTPUT_DIR}")
	23	+
	24	+if(CMAKE_SYSTEM_NAME STREQUAL "Windows")
	25	+ set(OPP_OS_TYPE "windows")
	26	+else()
	27	+ set(OPP_OS_TYPE "linux")
	28	+endif()
	29	+
	30	+string(TOLOWER "${CMAKE_SYSTEM_PROCESSOR}" CMAKE_SYSTEM_PROCESSOR_LOWER)
	31	+if(CMAKE_SYSTEM_PROCESSOR_LOWER MATCHES "^(aarch64\|arm64)$")
	32	+ set(OPP_CPU_TYPE "aarch64")
	33	+elseif(CMAKE_SYSTEM_PROCESSOR_LOWER MATCHES "^(x86_64\|amd64)$")
	34	+ set(OPP_CPU_TYPE "x86_64")
	35	+else()
	36	+ set(OPP_CPU_TYPE "${CMAKE_SYSTEM_PROCESSOR_LOWER}")
	37	+endif()
	38	+
	39	+set(CUSTOM_OP_OUTPUT_DIR "${PROJECT_OUTPUT_DIR}/op_graph/lib/${OPP_OS_TYPE}/${OPP_CPU_TYPE}")
	40	+file(MAKE_DIRECTORY "${CUSTOM_OP_OUTPUT_DIR}")
	41	+set(CUSTOM_OP_INCLUDE_DIR "${PROJECT_OUTPUT_DIR}/op_graph/include")
	42	+file(MAKE_DIRECTORY "${CUSTOM_OP_INCLUDE_DIR}")
	43	+
	44	+configure_file("${CMAKE_SOURCE_DIR}/ge/add_custom.h" "${CUSTOM_OP_INCLUDE_DIR}/add_custom.h" COPYONLY)
	45	+
	46	+set(KERNEL_SO_SOURCE_FILE "${CMAKE_SOURCE_DIR}/add_custom_kernel/add_kernel.so")
	47	+set(KERNEL_SO_OUTPUT_FILE "${CUSTOM_OP_OUTPUT_DIR}/add_kernel.so")
	48	+
	49	+set(ASCEND_HOME_PATH_OVERRIDE "" CACHE PATH "Optional ASCEND_HOME_PATH override")
	50	+if(ASCEND_HOME_PATH_OVERRIDE)
	51	+ set(ASCEND_HOME_PATH "${ASCEND_HOME_PATH_OVERRIDE}")
	52	+else()
	53	+ set(ASCEND_HOME_PATH "$ENV{ASCEND_HOME_PATH}")
	54	+endif()
	55	+
	56	+if(ASCEND_HOME_PATH)
	57	+ message(STATUS "ASCEND_HOME_PATH: ${ASCEND_HOME_PATH}")
	58	+else()
	59	+ message(WARNING "ASCEND_HOME_PATH is empty. Configure succeeds, but compilation requires a valid CANN toolkit path.")
	60	+endif()
	61	+
	62	+if(TILELANG_BUILD_CUSTOM_OP)
	63	+ add_library(cust_opapi SHARED
	64	+ ge/custom_op.cpp
	65	+ )
	66	+ target_compile_options(cust_opapi PRIVATE
	67	+ ${COMMON_COMPILE_OPTIONS}
	68	+ )
	69	+ target_compile_definitions(cust_opapi PRIVATE
	70	+ _GLIBCXX_USE_CXX11_ABI=0
	71	+ )
	72	+
	73	+ if(ASCEND_HOME_PATH)
	74	+ target_include_directories(cust_opapi PRIVATE
	75	+ "${CMAKE_SOURCE_DIR}/ge"
	76	+ "${ASCEND_HOME_PATH}/include"
	77	+ "${ASCEND_HOME_PATH}/include/graph"
	78	+ "${ASCEND_HOME_PATH}/include/register"
	79	+ "${ASCEND_HOME_PATH}/include/external"
	80	+ )
	81	+ target_link_directories(cust_opapi PRIVATE
	82	+ "${ASCEND_HOME_PATH}/lib64"
	83	+ )
	84	+ target_link_libraries(cust_opapi PRIVATE
	85	+ -Wl,--no-as-needed
	86	+ ascendcl
	87	+ lowering
	88	+ dl
	89	+ -Wl,--as-needed
	90	+ )
	91	+ endif()
	92	+
	93	+ set_target_properties(cust_opapi PROPERTIES
	94	+ OUTPUT_NAME "cust_opapi"
	95	+ LIBRARY_OUTPUT_DIRECTORY "${CUSTOM_OP_OUTPUT_DIR}"
	96	+ RUNTIME_OUTPUT_DIRECTORY "${CUSTOM_OP_OUTPUT_DIR}"
	97	+ ARCHIVE_OUTPUT_DIRECTORY "${CUSTOM_OP_OUTPUT_DIR}"
	98	+ )
	99	+
	100	+ add_custom_command(
	101	+ OUTPUT "${KERNEL_SO_OUTPUT_FILE}"
	102	+ COMMAND ${CMAKE_COMMAND} -E copy_if_different "${KERNEL_SO_SOURCE_FILE}" "${KERNEL_SO_OUTPUT_FILE}"
	103	+ DEPENDS "${KERNEL_SO_SOURCE_FILE}"
	104	+ COMMENT "Copying TileLang kernel .so to OPP package"
	105	+ )
	106	+ add_custom_target(tilelang_kernel_so ALL
	107	+ DEPENDS "${KERNEL_SO_OUTPUT_FILE}"
	108	+ )
	109	+
	110	+ install(FILES "${KERNEL_SO_OUTPUT_FILE}"
	111	+ DESTINATION "${CUSTOM_OP_OUTPUT_DIR}"
	112	+ )
	113	+ install(FILES "${CUSTOM_OP_INCLUDE_DIR}/add_custom.h"
	114	+ DESTINATION "${CUSTOM_OP_INCLUDE_DIR}"
	115	+ )
	116	+endif()
	117	+
	118	+if(TILELANG_BUILD_SESSION_RUN)
	119	+ add_executable(tilelang_session_run
	120	+ session_run/main.cc
	121	+ )
	122	+ target_compile_options(tilelang_session_run PRIVATE ${COMMON_COMPILE_OPTIONS})
	123	+ target_compile_definitions(tilelang_session_run PRIVATE
	124	+ _GLIBCXX_USE_CXX11_ABI=0
	125	+ )
	126	+
	127	+ if(ASCEND_HOME_PATH)
	128	+ target_include_directories(tilelang_session_run PRIVATE
	129	+ "${CUSTOM_OP_INCLUDE_DIR}"
	130	+ "${CMAKE_SOURCE_DIR}/ge"
	131	+ "${ASCEND_HOME_PATH}/include"
	132	+ "${ASCEND_HOME_PATH}/include/graph"
	133	+ "${ASCEND_HOME_PATH}/include/ge"
	134	+ "${ASCEND_HOME_PATH}/opp/built-in/op_proto/inc"
	135	+ )
	136	+ target_link_directories(tilelang_session_run PRIVATE "${ASCEND_HOME_PATH}/lib64")
	137	+ target_link_libraries(tilelang_session_run PRIVATE
	138	+ graph
	139	+ ge_runner
	140	+ ge_compiler
	141	+ ascendcl
	142	+ graph_base
	143	+ )
	144	+ endif()
	145	+
	146	+ set_target_properties(tilelang_session_run PROPERTIES
	147	+ RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}"
	148	+ )
	149	+endif()

Aexamples/custom_op/tilelang_add_custom/README.md+177-0

@@ -0,0 +1,177 @@
	1	+# TileLang Add Custom Operator 样例
	2	+
	3	+## 样例概述
	4	+
	5	+- 构图入口: GE 原生 (Session API)
	6	+- 算子编程语言: TileLang
	7	+- 编译方式: TileLang 预编译产出 host-wrapper `.so`，GE 运行时通过 `dlopen` 加载
	8	+- 核心链路: `TileLang kernel → 预编译 .so → GE 交付件 → 进程内构图 → Session::ExecuteGraphWithStreamAsync 在线执行`
	9	+- 场景: 场景 A — 动态图在线执行（预编译 kernel + host 调度）
	10	+
	11	+本样例以 element-wise Add 算子为例，展示如何将 TileLang 编写的 kernel 通过 GE 语言无关自定义算子机制接入图编译和执行流程。
	12	+
	13	+## 目录结构
	14	+
	15	+```text
	16	+tilelang_add_custom/
	17	+├── README.md
	18	+├── README_en.md
	19	+├── CMakeLists.txt # 构建 libcust_opapi.so + session_run + 安装 add_kernel.so
	20	+├── run.sh # 一键编译运行
	21	+├── add_custom_kernel/
	22	+│ └── add_custom_kernel.py # TileLang kernel + 编译产出 add_kernel.so
	23	+├── ge/
	24	+│ ├── add_custom.h # REG_OP 算子 proto 定义
	25	+│ └── custom_op.cpp # EagerExecuteOp + ShapeInferOp 实现
	26	+└── session_run/
	27	+ └── main.cc # GE 原生构图 + Session 执行 + 精度校验
	28	+```
	29	+
	30	+## 核心流程
	31	+
	32	+```text
	33	+TileLang kernel 源码 (add_custom_kernel.py)
	34	+ ↓ TileLang-Ascend 编译器 (TVM + Ascend C codegen + Bisheng)
	35	+add_kernel.so (host-wrapper，导出 call 函数)
	36	+ ↓ dlopen + dlsym("call")
	37	+GE 自定义算子 (AddCustom, EagerExecuteOp)
	38	+ ↓ call(x_ptr, y_ptr, z_ptr, stream) — 内部封装 main_kernel<<<>>> launch
	39	+NPU 执行
	40	+```
	41	+
	42	+TileLang-Ascend 编译后的 `.so` 导出函数签名为：
	43	+
	44	+```c
	45	+extern "C" void call(uint8_t* A_handle, uint8_t* B_handle, uint8_t* C_handle, aclrtStream stream)
	46	+```
	47	+
	48	+该函数内部封装了 `main_kernel<<<>>>` 的 launch 逻辑（含硬件调度地址获取、tiling 等），GE 侧无需手动拼装 args。
	49	+
	50	+## 前置依赖
	51	+
	52	+### CANN
	53	+
	54	+- 已正确安装并配置 CANN 环境（`source ${ASCEND_HOME_PATH}/set_env.sh`）
	55	+- 当前环境具备 ACL、GE、Graph 相关头文件与库
	56	+
	57	+### TileLang-Ascend
	58	+
	59	+需安装 TileLang 主包和 TileLang-Ascend 后端：
	60	+
	61	+```bash
	62	+pip install tilelang # 主包
	63	+# TileLang-Ascend 后端：从 https://github.com/tile-ai/tilelang-ascend 安装
	64	+```
	65	+
	66	+若 TileLang-Ascend 以源码方式安装（未 `pip install`），需设置环境变量：
	67	+
	68	+```bash
	69	+export TILELANG_ASCEND_HOME=/path/to/tilelang-ascend
	70	+```
	71	+
	72	+### 环境变量
	73	+
	74	+\| 变量 \| 必需 \| 说明 \|
	75	+\|------\|------\|------\|
	76	+\| `ASCEND_HOME_PATH` \| 是 \| CANN toolkit 路径 \|
	77	+\| `TILELANG_ASCEND_HOME` \| 否 \| TileLang-Ascend 源码安装路径（pip 安装则无需设置） \|
	78	+\| `ASCEND_CUSTOM_OPP_PATH` \| 自动 \| 由 `run.sh` 自动设置 \|
	79	+
	80	+## 快速运行
	81	+
	82	+```bash
	83	+source ${ASCEND_HOME_PATH}/set_env.sh
	84	+bash run.sh
	85	+```
	86	+
	87	+`run.sh` 依次执行 4 个步骤：
	88	+
	89	+1. 编译 TileLang kernel，产出 `add_kernel.so`
	90	+2. 构建 `libcust_opapi.so` 和 `tilelang_session_run`，将 `add_kernel.so` 安装到 OPP 包
	91	+3. 确认 kernel `.so` 已在 OPP 包中
	92	+4. 运行测试程序
	93	+
	94	+成功时终端输出：
	95	+
	96	+```text
	97	+[INFO] Step 1/4: compile TileLang kernel
	98	+Kernel .so saved to: add_kernel.so
	99	+[INFO] Step 2/4: build custom op library and session_run
	100	+...
	101	+[INFO] Step 3/4: kernel .so installed in OPP package.
	102	+[INFO] Step 4/4: run session test
	103	+Precision check passed, max_error=0
	104	+[INFO] Sample pipeline finished.
	105	+```
	106	+
	107	+## 关键文件说明
	108	+
	109	+### `ge/custom_op.cpp`
	110	+
	111	+GE 交付件，实现 `EagerExecuteOp` + `ShapeInferOp`：
	112	+
	113	+- Execute:
	114	+ 1. 首次调用时通过 `dlopen` 加载 `add_kernel.so`（路径从 `ASCEND_CUSTOM_OPP_PATH` 定位），`dlsym` 获取 `call` 函数指针
	115	+ 2. 校验两个输入的 shape size 均为 4096
	116	+ 3. 分配输出 Tensor，调用 `call(x_ptr, y_ptr, z_ptr, stream)`
	117	+- InferShape / InferDataType: 输出 shape 和 dtype 与输入相同
	118	+- 使用 `std::once_flag` 保证线程安全的延迟加载
	119	+- kernel `.so` 路径从 `ASCEND_CUSTOM_OPP_PATH` 环境变量定位，不依赖工作目录
	120	+
	121	+### `ge/add_custom.h`
	122	+
	123	+`REG_OP(AddCustom)` 声明算子的输入输出规格，供 GE 原生构图创建节点。
	124	+
	125	+### `add_custom_kernel/add_custom_kernel.py`
	126	+
	127	+TileLang kernel 源码，定义 element-wise Add 并编译产出 `add_kernel.so`。
	128	+
	129	+### `session_run/main.cc`
	130	+
	131	+GE 原生构图测试程序：
	132	+
	133	+1. `GEInitialize` + 创建 `Session`
	134	+2. 构建 `Data → AddCustom` 计算图
	135	+3. `AddGraph` → `CompileGraph` → `LoadGraph`
	136	+4. 分配 device 内存，H2D 拷贝输入数据
	137	+5. `ExecuteGraphWithStreamAsync` 执行
	138	+6. D2H 拷贝输出，逐元素精度校验（含 NaN 检查）
	139	+
	140	+## 算子规格
	141	+
	142	+\| 项目 \| 内容 \|
	143	+\|------\|------\|
	144	+\| 算子类型 \| `AddCustom` \|
	145	+\| 输入 \| `x` (float32), `y` (float32) \|
	146	+\| 输出 \| `z` (float32) \|
	147	+\| 输入 shape \| `[4096]` (固定) \|
	148	+\| 输出 shape \| `[4096]` \|
	149	+\| 格式 \| ND \|
	150	+\| kernel 名称 \| `main_kernel`（由 `call` 封装） \|
	151	+\| BLOCK_SIZE \| 1024 \|
	152	+
	153	+## 分步运行
	154	+
	155	+```bash
	156	+# 1. 编译 TileLang kernel
	157	+cd add_custom_kernel && python3 add_custom_kernel.py && cd ..
	158	+
	159	+# 2. 构建（含将 add_kernel.so 安装到 OPP 包）
	160	+cmake -S . -B build -DCMAKE_BUILD_TYPE=Release
	161	+cmake --build build -j$(nproc)
	162	+cmake --install build
	163	+
	164	+# 3. 配置环境变量
	165	+export ASCEND_CUSTOM_OPP_PATH="$(pwd)/output:$ASCEND_CUSTOM_OPP_PATH"
	166	+
	167	+# 4. 运行
	168	+./build/tilelang_session_run
	169	+```
	170	+
	171	+## 注意事项
	172	+
	173	+- kernel 编译时固定 N=4096，Execute 中会校验输入 shape size 是否匹配，不匹配则返回失败。
	174	+- `ge.graphRunMode=1` 确保走在线执行链路（PRIORITY_GRAPH 模式）。
	175	+- 当前样例仅支持 float32，如需支持更多数据类型需调整 `REG_OP` 的 `DATATYPE` 约束和 TileLang kernel 的 dtype 参数。
	176	+- TileLang-Ascend 的平台检测基于 `torch.npu.get_device_name()`，Ascend910 映射为 A2 平台。
	177	+- kernel `.so` 安装在 OPP 包的 `op_graph/lib/<os>/<arch>/` 目录下，与 `libcust_opapi.so` 同目录，路径从 `ASCEND_CUSTOM_OPP_PATH` 定位。

Aexamples/custom_op/tilelang_add_custom/README_en.md+86-0

@@ -0,0 +1,86 @@
	1	+# TileLang Add Custom Operator Sample
	2	+
	3	+## Sample Overview
	4	+
	5	+- Graph construction entry: GE native (Session API)
	6	+- Operator programming language: TileLang
	7	+- Compilation method: TileLang pre-compiles to host-wrapper `.so`, loaded via `dlopen` at runtime
	8	+- Core pipeline: `TileLang kernel → pre-compiled .so → GE deliverable → in-process graph → Session::ExecuteGraphWithStreamAsync online execution`
	9	+- Scenario: Scenario A — online execution with pre-compiled kernel
	10	+
	11	+This sample demonstrates how to integrate a TileLang kernel into GE's graph compilation and execution flow via the language-independent custom operator mechanism, using an element-wise Add operator as an example.
	12	+
	13	+## Directory Structure
	14	+
	15	+```text
	16	+tilelang_add_custom/
	17	+├── README.md
	18	+├── README_en.md
	19	+├── CMakeLists.txt # Build libcust_opapi.so + session_run + install add_kernel.so
	20	+├── run.sh # One-click build and run
	21	+├── add_custom_kernel/
	22	+│ └── add_custom_kernel.py # TileLang kernel + compile to add_kernel.so
	23	+├── ge/
	24	+│ ├── add_custom.h # REG_OP proto definition
	25	+│ └── custom_op.cpp # EagerExecuteOp + ShapeInferOp implementation
	26	+└── session_run/
	27	+ └── main.cc # GE native graph + Session execution + precision check
	28	+```
	29	+
	30	+## Core Pipeline
	31	+
	32	+```text
	33	+TileLang kernel source (add_custom_kernel.py)
	34	+ ↓ TileLang-Ascend compiler (TVM + Ascend C codegen + Bisheng)
	35	+add_kernel.so (host-wrapper, exports call function)
	36	+ ↓ dlopen + dlsym("call")
	37	+GE custom operator (AddCustom, EagerExecuteOp)
	38	+ ↓ call(x_ptr, y_ptr, z_ptr, stream) — wraps main_kernel<<<>>> launch
	39	+NPU execution
	40	+```
	41	+
	42	+## Prerequisites
	43	+
	44	+### CANN
	45	+
	46	+- CANN environment properly installed and configured (`source ${ASCEND_HOME_PATH}/set_env.sh`)
	47	+
	48	+### TileLang-Ascend
	49	+
	50	+```bash
	51	+pip install tilelang
	52	+# TileLang-Ascend backend: install from https://github.com/tile-ai/tilelang-ascend
	53	+```
	54	+
	55	+If TileLang-Ascend is installed from source (not via pip), set:
	56	+
	57	+```bash
	58	+export TILELANG_ASCEND_HOME=/path/to/tilelang-ascend
	59	+```
	60	+
	61	+## Quick Start
	62	+
	63	+```bash
	64	+source ${ASCEND_HOME_PATH}/set_env.sh
	65	+bash run.sh
	66	+```
	67	+
	68	+## Operator Specification
	69	+
	70	+\| Item \| Value \|
	71	+\|------\|-------\|
	72	+\| Op type \| `AddCustom` \|
	73	+\| Inputs \| `x` (float32), `y` (float32) \|
	74	+\| Output \| `z` (float32) \|
	75	+\| Input shape \| `[4096]` (fixed) \|
	76	+\| Format \| ND \|
	77	+\| Kernel name \| `main_kernel` (wrapped by `call`) \|
	78	+\| BLOCK_SIZE \| 1024 \|
	79	+
	80	+## Notes
	81	+
	82	+- The kernel is compiled with fixed N=4096. Execute validates input shape size and returns failure on mismatch.
	83	+- `ge.graphRunMode=1` ensures online execution (PRIORITY_GRAPH mode).
	84	+- Only float32 is supported. To support more data types, adjust `REG_OP` DATATYPE and TileLang kernel dtype parameter.
	85	+- TileLang-Ascend platform detection is based on `torch.npu.get_device_name()`. Ascend910 maps to A2 platform.
	86	+- The kernel `.so` is installed in the OPP package at `op_graph/lib/<os>/<arch>/`, alongside `libcust_opapi.so`.

Aexamples/custom_op/tilelang_add_custom/add_custom_kernel/add_custom_kernel.py+69-0

@@ -0,0 +1,69 @@
	1	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	2	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	3	+# CANN Open Software License Agreement Version 2.0 (the "License").
	4	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	5	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	6	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	7	+# See LICENSE in the root of the software repository for the full text of the License.
	8	+
	9	+"""
	10	+TileLang Add Kernel + 编译产出 .so
	11	+
	12	+本文件用 TileLang 实现 element-wise Add kernel，并编译产出 Ascend .so 交付件。
	13	+产出的 .so 供 GE 自定义算子 (EagerExecuteOp) 通过 dlopen + dlsym("call") 加载并调用。
	14	+
	15	+TileLang-Ascend 编译后的 .so 导出函数签名为:
	16	+ extern "C" void call(uint8_t* A_handle, uint8_t* B_handle, uint8_t* C_handle, aclrtStream stream)
	17	+内部封装了 main_kernel<<<>>> 的 launch 逻辑。
	18	+"""
	19	+
	20	+import os
	21	+import shutil
	22	+
	23	+import tilelang
	24	+import tilelang.language as T
	25	+
	26	+N = 4096
	27	+BLOCK_SIZE = 1024
	28	+
	29	+
	30	+@tilelang.jit(out_idx=[-1])
	31	+def vec_add(n, block_size, dtype="float"):
	32	+ m_num = n // block_size
	33	+ vec_num = 2
	34	+
	35	+ @T.prim_func
	36	+ def main(
	37	+ a: T.Tensor((n,), dtype),
	38	+ b: T.Tensor((n,), dtype),
	39	+ c: T.Tensor((n,), dtype),
	40	+ ):
	41	+ with T.Kernel(m_num, is_npu=True) as (cid, vid):
	42	+ a_ub = T.alloc_ub((block_size // vec_num,), dtype)
	43	+ b_ub = T.alloc_ub((block_size // vec_num,), dtype)
	44	+ c_ub = T.alloc_ub((block_size // vec_num,), dtype)
	45	+ with T.Scope("V"):
	46	+ T.copy(a[cid * block_size + vid * block_size // vec_num], a_ub)
	47	+ T.copy(b[cid * block_size + vid * block_size // vec_num], b_ub)
	48	+
	49	+ T.barrier_all()
	50	+ T.tile.add(c_ub, a_ub, b_ub)
	51	+ T.barrier_all()
	52	+
	53	+ T.copy(c_ub, c[cid * block_size + vid * block_size // vec_num])
	54	+
	55	+ return main
	56	+
	57	+
	58	+func = vec_add(N, BLOCK_SIZE)
	59	+
	60	+adapter = func.adapter
	61	+so_path = adapter.lib._name
	62	+
	63	+output_dir = os.path.dirname(os.path.abspath(__file__))
	64	+output_path = os.path.join(output_dir, "add_kernel.so")
	65	+shutil.copy2(so_path, output_path)
	66	+
	67	+print(f"Kernel .so saved to: {output_path}")
	68	+print(f"File size: {os.path.getsize(output_path)} bytes")
	69	+print("Export function: call(uint8_t* A, uint8_t* B, uint8_t* C, aclrtStream stream)")

Aexamples/custom_op/tilelang_add_custom/ge/add_custom.h+25-0

@@ -0,0 +1,25 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#ifndef EXAMPLES_CUSTOM_OP_TILELANG_ADD_CUSTOM_GE_ADD_CUSTOM_H_
	12	+#define EXAMPLES_CUSTOM_OP_TILELANG_ADD_CUSTOM_GE_ADD_CUSTOM_H_
	13	+
	14	+#include "graph/operator_reg.h"
	15	+
	16	+namespace ge {
	17	+REG_OP(AddCustom)
	18	+ .INPUT(x, "T")
	19	+ .INPUT(y, "T")
	20	+ .OUTPUT(z, "T")
	21	+ .DATATYPE(T, TensorType({DT_FLOAT}))
	22	+ .OP_END_FACTORY_REG(AddCustom);
	23	+} // namespace ge
	24	+
	25	+#endif // EXAMPLES_CUSTOM_OP_TILELANG_ADD_CUSTOM_GE_ADD_CUSTOM_H_

Aexamples/custom_op/tilelang_add_custom/ge/custom_op.cpp+135-0

@@ -0,0 +1,135 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#include <dlfcn.h>
	12	+#include <iostream>
	13	+#include <mutex>
	14	+#include <string>
	15	+#include "graph/custom_op.h"
	16	+#include "acl/acl_rt.h"
	17	+
	18	+using namespace ge;
	19	+
	20	+namespace {
	21	+constexpr const char *kKernelSoName = "add_kernel.so";
	22	+constexpr const char *kCallFuncName = "call";
	23	+constexpr int64_t kExpectedNumElements = 4096;
	24	+
	25	+using CallFunc = void ()(void x_ptr, void y_ptr, void z_ptr, void *stream);
	26	+
	27	+std::string GetKernelSoPath() {
	28	+ const char *opp_path = std::getenv("ASCEND_CUSTOM_OPP_PATH");
	29	+ if (opp_path == nullptr \|\| opp_path[0] == '\0') {
	30	+ return std::string(kKernelSoName);
	31	+ }
	32	+ std::string path(opp_path);
	33	+ size_t colon = path.find(':');
	34	+ if (colon != std::string::npos) {
	35	+ path = path.substr(0, colon);
	36	+ }
	37	+ if (!path.empty() && path.back() != '/') {
	38	+ path += '/';
	39	+ }
	40	+ path += "op_graph/lib/linux/aarch64/";
	41	+ path += kKernelSoName;
	42	+ return path;
	43	+}
	44	+} // namespace
	45	+
	46	+class AddCustom : public EagerExecuteOp, public ShapeInferOp {
	47	+ public:
	48	+ ~AddCustom() {
	49	+ if (so_handle_ != nullptr) {
	50	+ (void)dlclose(so_handle_);
	51	+ }
	52	+ }
	53	+
	54	+ graphStatus Execute(gert::EagerOpExecutionContext *ctx) override {
	55	+ {
	56	+ std::call_once(load_flag_, [this]() { load_status_ = LoadKernel(); });
	57	+ }
	58	+ if (load_status_ != GRAPH_SUCCESS) {
	59	+ return GRAPH_FAILED;
	60	+ }
	61	+
	62	+ const gert::Tensor *input_x = ctx->GetInputTensor(0);
	63	+ const gert::Tensor *input_y = ctx->GetInputTensor(1);
	64	+ if (input_x == nullptr \|\| input_y == nullptr) {
	65	+ std::cerr << "GetInputTensor failed" << std::endl;
	66	+ return GRAPH_FAILED;
	67	+ }
	68	+
	69	+ int64_t x_size = input_x->GetShapeSize();
	70	+ int64_t y_size = input_y->GetShapeSize();
	71	+ if (x_size != kExpectedNumElements \|\| y_size != kExpectedNumElements) {
	72	+ std::cerr << "Input shape size mismatch: x=" << x_size << ", y=" << y_size
	73	+ << ", expected=" << kExpectedNumElements << std::endl;
	74	+ return GRAPH_FAILED;
	75	+ }
	76	+
	77	+ gert::Tensor *output_z =
	78	+ ctx->MallocOutputTensor(0, input_x->GetShape(), input_x->GetFormat(), input_x->GetDataType());
	79	+ if (output_z == nullptr) {
	80	+ std::cerr << "MallocOutputTensor failed" << std::endl;
	81	+ return GRAPH_FAILED;
	82	+ }
	83	+
	84	+ void *stream = ctx->GetStream();
	85	+ call_func_(const_cast<void >(input_x->GetAddr()), const_cast<void >(input_y->GetAddr()), output_z->GetAddr(),
	86	+ stream);
	87	+ return GRAPH_SUCCESS;
	88	+ }
	89	+
	90	+ graphStatus InferShape(gert::InferShapeContext *ctx) override {
	91	+ const auto *input_shape = ctx->GetInputShape(0);
	92	+ auto *output_shape = ctx->GetOutputShape(0);
	93	+ if (input_shape == nullptr \|\| output_shape == nullptr) {
	94	+ return GRAPH_FAILED;
	95	+ }
	96	+ output_shape->SetDimNum(input_shape->GetDimNum());
	97	+ for (size_t i = 0; i < input_shape->GetDimNum(); ++i) {
	98	+ output_shape->SetDim(i, input_shape->GetDim(i));
	99	+ }
	100	+ return GRAPH_SUCCESS;
	101	+ }
	102	+
	103	+ graphStatus InferDataType(gert::InferDataTypeContext *ctx) override {
	104	+ return ctx->SetOutputDataType(0, ctx->GetInputDataType(0));
	105	+ }
	106	+
	107	+ private:
	108	+ graphStatus LoadKernel() {
	109	+ std::string so_path = GetKernelSoPath();
	110	+ so_handle_ = dlopen(so_path.c_str(), RTLD_NOW);
	111	+ if (so_handle_ == nullptr) {
	112	+ std::cerr << "dlopen failed: " << dlerror() << std::endl;
	113	+ return GRAPH_FAILED;
	114	+ }
	115	+
	116	+ dlerror();
	117	+ call_func_ = reinterpret_cast<CallFunc>(dlsym(so_handle_, kCallFuncName));
	118	+ const char *error = dlerror();
	119	+ if (error != nullptr) {
	120	+ std::cerr << "dlsym '" << kCallFuncName << "' failed: " << error << std::endl;
	121	+ (void)dlclose(so_handle_);
	122	+ so_handle_ = nullptr;
	123	+ return GRAPH_FAILED;
	124	+ }
	125	+
	126	+ return GRAPH_SUCCESS;
	127	+ }
	128	+
	129	+ std::once_flag load_flag_;
	130	+ graphStatus load_status_ = GRAPH_FAILED;
	131	+ void *so_handle_ = nullptr;
	132	+ CallFunc call_func_ = nullptr;
	133	+};
	134	+
	135	+REG_AUTO_MAPPING_OP(AddCustom);

Aexamples/custom_op/tilelang_add_custom/run.sh+79-0

@@ -0,0 +1,79 @@
	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	+PROJECT_DIR="${SCRIPT_DIR}"
	16	+BUILD_DIR="${PROJECT_DIR}/build"
	17	+OUTPUT_DIR="${PROJECT_DIR}/output"
	18	+
	19	+info() {
	20	+ echo "[INFO] $*"
	21	+}
	22	+
	23	+error() {
	24	+ echo "[ERROR] $*" >&2
	25	+}
	26	+
	27	+if [[ -z "${ASCEND_HOME_PATH:-}" ]]; then
	28	+ error "ASCEND_HOME_PATH is empty. Please source CANN set_env.sh first."
	29	+ exit 1
	30	+fi
	31	+
	32	+if [[ -n "${TILELANG_ASCEND_HOME:-}" ]]; then
	33	+ export PYTHONPATH="${TILELANG_ASCEND_HOME}:${PYTHONPATH:-}"
	34	+ export LD_LIBRARY_PATH="${TILELANG_ASCEND_HOME}/build:${LD_LIBRARY_PATH:-}"
	35	+ info "Using TILELANG_ASCEND_HOME=${TILELANG_ASCEND_HOME}"
	36	+else
	37	+ if ! python3 -c "import tilelang" 2>/dev/null; then
	38	+ error "tilelang is not importable. Set TILELANG_ASCEND_HOME or install tilelang-ascend."
	39	+ exit 1
	40	+ fi
	41	+fi
	42	+
	43	+mkdir -p "${BUILD_DIR}" "${OUTPUT_DIR}"
	44	+
	45	+# Step 1: 编译 TileLang kernel，产出 .so
	46	+info "Step 1/4: compile TileLang kernel"
	47	+(
	48	+ cd "${PROJECT_DIR}/add_custom_kernel"
	49	+ python3 add_custom_kernel.py
	50	+)
	51	+if [[ ! -f "${PROJECT_DIR}/add_custom_kernel/add_kernel.so" ]]; then
	52	+ error "TileLang kernel .so not found: add_custom_kernel/add_kernel.so"
	53	+ error "Please check TileLang installation and kernel compilation output."
	54	+ exit 1
	55	+fi
	56	+info "TileLang kernel .so generated."
	57	+
	58	+# Step 2: 构建 libcust_opapi.so + session_run（含将 add_kernel.so 安装到 OPP 包）
	59	+info "Step 2/4: build custom op library and session_run"
	60	+cmake -S "${PROJECT_DIR}" -B "${BUILD_DIR}" -DCMAKE_BUILD_TYPE=Release
	61	+cmake --build "${BUILD_DIR}" -j"$(nproc 2>/dev/null \|\| echo 8)"
	62	+cmake --install "${BUILD_DIR}"
	63	+
	64	+export ASCEND_CUSTOM_OPP_PATH="${OUTPUT_DIR}:${ASCEND_CUSTOM_OPP_PATH:-}"
	65	+info "ASCEND_CUSTOM_OPP_PATH=${ASCEND_CUSTOM_OPP_PATH}"
	66	+
	67	+# Step 3: 确认 kernel .so 已在 OPP 包中
	68	+KERNEL_SO="${OUTPUT_DIR}/op_graph/lib/linux/$(uname -m \| tr '[:upper:]' '[:lower:]' \| sed 's/x86_64/x86_64/;s/aarch64/aarch64/')/add_kernel.so"
	69	+if [[ ! -f "${KERNEL_SO}" ]]; then
	70	+ error "kernel .so not found in OPP package: ${KERNEL_SO}"
	71	+ exit 1
	72	+fi
	73	+info "kernel .so installed in OPP package."
	74	+
	75	+# Step 4: 运行测试
	76	+info "Step 4/4: run session test"
	77	+"${BUILD_DIR}/tilelang_session_run"
	78	+
	79	+info "Sample pipeline finished."

Aexamples/custom_op/tilelang_add_custom/session_run/main.cc+204-0

@@ -0,0 +1,204 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#include <cmath>
	12	+#include <cstdlib>
	13	+#include <map>
	14	+#include <memory>
	15	+#include <random>
	16	+#include <vector>
	17	+
	18	+#include "acl/acl_rt.h"
	19	+#include "ge/ge_api.h"
	20	+#include "graph.h"
	21	+#include "ops_proto_legacy.h"
	22	+#include "tensor.h"
	23	+#include "types.h"
	24	+#include "add_custom.h"
	25	+
	26	+using ge::Operator;
	27	+
	28	+namespace {
	29	+constexpr uint32_t kGraphId = 0;
	30	+constexpr int64_t kNumElements = 4096;
	31	+constexpr size_t kDataSizeBytes = static_cast<size_t>(kNumElements) * sizeof(float);
	32	+constexpr int kRandomSeed = 42;
	33	+constexpr float kErrorTolerance = 1e-5f;
	34	+constexpr int kMaxErrorDetails = 10;
	35	+constexpr size_t kNumInputs = 2U;
	36	+
	37	+#define CHECK_ACL(ret, msg) \
	38	+ do { \
	39	+ if ((ret) != ACL_ERROR_NONE) { \
	40	+ std::cerr << (msg) << ", aclError: " << (ret) << std::endl; \
	41	+ return 1; \
	42	+ } \
	43	+ } while (0)
	44	+
	45	+std::unique_ptr<ge::Graph> BuildGraph() {
	46	+ ge::TensorDesc input_desc(ge::Shape({kNumElements}), ge::FORMAT_ND, ge::DT_FLOAT);
	47	+
	48	+ auto data_x = ge::op::Data("data_x");
	49	+ data_x.update_input_desc_x(input_desc);
	50	+ data_x.update_output_desc_y(input_desc);
	51	+ auto data_y = ge::op::Data("data_y");
	52	+ data_y.update_input_desc_x(input_desc);
	53	+ data_y.update_output_desc_y(input_desc);
	54	+
	55	+ auto add = ge::op::AddCustom("add").set_input_x(data_x).set_input_y(data_y);
	56	+
	57	+ std::vector<Operator> inputs = {data_x, data_y};
	58	+ std::vector<Operator> outputs = {add};
	59	+
	60	+ auto graph = std::make_unique<ge::Graph>("tilelang_add_graph");
	61	+ graph->SetInputs(inputs).SetOutputs(outputs);
	62	+ return graph;
	63	+}
	64	+
	65	+bool VerifyResult(const std::vector<float> &host_x, const std::vector<float> &host_y,
	66	+ const std::vector<float> &host_z) {
	67	+ int error_count = 0;
	68	+ float max_error = 0.0f;
	69	+ for (int64_t i = 0; i < kNumElements; ++i) {
	70	+ float expected = host_x[i] + host_y[i];
	71	+ if (std::isnan(host_z[i]) \|\| std::isnan(expected)) {
	72	+ std::cerr << "NaN detected at [" << i << "]: got=" << host_z[i] << ", expected=" << expected << std::endl;
	73	+ error_count++;
	74	+ continue;
	75	+ }
	76	+ float error = std::abs(host_z[i] - expected);
	77	+ max_error = std::max(max_error, error);
	78	+ if (error > kErrorTolerance) {
	79	+ if (error_count < kMaxErrorDetails) {
	80	+ std::cerr << "Error at [" << i << "]: expected=" << expected << ", got=" << host_z[i] << std::endl;
	81	+ }
	82	+ error_count++;
	83	+ }
	84	+ }
	85	+ if (error_count > 0) {
	86	+ std::cerr << "Precision check failed: " << error_count << " errors, max_error=" << max_error << std::endl;
	87	+ return false;
	88	+ }
	89	+ std::cout << "Precision check passed, max_error=" << max_error << std::endl;
	90	+ return true;
	91	+}
	92	+} // namespace
	93	+
	94	+int main(int argc, char *argv[]) {
	95	+ (void)argc;
	96	+ (void)argv;
	97	+
	98	+ std::map<ge::AscendString, ge::AscendString> options = {
	99	+ {"ge.exec.deviceId", "0"},
	100	+ {"ge.graphRunMode", "1"},
	101	+ };
	102	+
	103	+ auto init_ret = ge::GEInitialize(options);
	104	+ if (init_ret != ge::SUCCESS) {
	105	+ std::cerr << "GEInitialize failed, ret: " << init_ret << std::endl;
	106	+ return 1;
	107	+ }
	108	+
	109	+ aclrtStream stream = nullptr;
	110	+ CHECK_ACL(aclrtCreateStream(&stream), "Failed to create stream");
	111	+
	112	+ int ret_code = 0;
	113	+ {
	114	+ ge::Session session(options);
	115	+ auto graph = BuildGraph();
	116	+
	117	+ auto ret = session.AddGraph(kGraphId, *graph);
	118	+ if (ret != ge::SUCCESS) {
	119	+ std::cerr << "AddGraph failed, ret: " << ret << std::endl;
	120	+ ret_code = 1;
	121	+ }
	122	+
	123	+ if (ret_code == 0) {
	124	+ ret = session.CompileGraph(kGraphId);
	125	+ if (ret != ge::SUCCESS) {
	126	+ std::cerr << "CompileGraph failed, ret: " << ret << std::endl;
	127	+ ret_code = 1;
	128	+ }
	129	+ }
	130	+
	131	+ if (ret_code == 0) {
	132	+ std::map<ge::AscendString, ge::AscendString> load_options;
	133	+ ret = session.LoadGraph(kGraphId, load_options, stream);
	134	+ if (ret != ge::SUCCESS) {
	135	+ std::cerr << "LoadGraph failed, ret: " << ret << std::endl;
	136	+ ret_code = 1;
	137	+ }
	138	+ }
	139	+
	140	+ if (ret_code == 0) {
	141	+ void *x_ptr = nullptr;
	142	+ void *y_ptr = nullptr;
	143	+ void *z_ptr = nullptr;
	144	+ CHECK_ACL(aclrtMalloc(&x_ptr, kDataSizeBytes, ACL_MEM_MALLOC_HUGE_FIRST), "aclrtMalloc x failed");
	145	+ CHECK_ACL(aclrtMalloc(&y_ptr, kDataSizeBytes, ACL_MEM_MALLOC_HUGE_FIRST), "aclrtMalloc y failed");
	146	+ CHECK_ACL(aclrtMalloc(&z_ptr, kDataSizeBytes, ACL_MEM_MALLOC_HUGE_FIRST), "aclrtMalloc z failed");
	147	+
	148	+ std::vector<float> host_x(kNumElements);
	149	+ std::vector<float> host_y(kNumElements);
	150	+ std::vector<float> host_z(kNumElements);
	151	+ std::mt19937 rng(kRandomSeed);
	152	+ std::uniform_real_distribution<float> dist(0.0f, 1.0f);
	153	+ for (int64_t i = 0; i < kNumElements; ++i) {
	154	+ host_x[i] = dist(rng);
	155	+ host_y[i] = dist(rng);
	156	+ }
	157	+ CHECK_ACL(aclrtMemcpy(x_ptr, kDataSizeBytes, host_x.data(), kDataSizeBytes, ACL_MEMCPY_HOST_TO_DEVICE),
	158	+ "aclrtMemcpy x H2D failed");
	159	+ CHECK_ACL(aclrtMemcpy(y_ptr, kDataSizeBytes, host_y.data(), kDataSizeBytes, ACL_MEMCPY_HOST_TO_DEVICE),
	160	+ "aclrtMemcpy y H2D failed");
	161	+
	162	+ std::vector<gert::Tensor> inputs(kNumInputs);
	163	+ inputs[0] = {{{kNumElements}, {kNumElements}},
	164	+ {ge::FORMAT_ND, ge::FORMAT_ND, {}},
	165	+ gert::kOnDeviceHbm,
	166	+ ge::DT_FLOAT,
	167	+ x_ptr};
	168	+ inputs[1] = {{{kNumElements}, {kNumElements}},
	169	+ {ge::FORMAT_ND, ge::FORMAT_ND, {}},
	170	+ gert::kOnDeviceHbm,
	171	+ ge::DT_FLOAT,
	172	+ y_ptr};
	173	+
	174	+ std::vector<gert::Tensor> outputs(1);
	175	+ outputs[0] = {{{kNumElements}, {kNumElements}},
	176	+ {ge::FORMAT_ND, ge::FORMAT_ND, {}},
	177	+ gert::kOnDeviceHbm,
	178	+ ge::DT_FLOAT,
	179	+ z_ptr};
	180	+
	181	+ ret = session.ExecuteGraphWithStreamAsync(kGraphId, stream, inputs, outputs);
	182	+ if (ret != ge::SUCCESS) {
	183	+ std::cerr << "ExecuteGraphWithStreamAsync failed, ret: " << ret << std::endl;
	184	+ ret_code = 1;
	185	+ } else {
	186	+ CHECK_ACL(aclrtSynchronizeStream(stream), "aclrtSynchronizeStream failed");
	187	+ CHECK_ACL(aclrtMemcpy(host_z.data(), kDataSizeBytes, z_ptr, kDataSizeBytes, ACL_MEMCPY_DEVICE_TO_HOST),
	188	+ "aclrtMemcpy z D2H failed");
	189	+ if (!VerifyResult(host_x, host_y, host_z)) {
	190	+ ret_code = 1;
	191	+ }
	192	+ }
	193	+ CHECK_ACL(aclrtFree(x_ptr), "aclrtFree x failed");
	194	+ CHECK_ACL(aclrtFree(y_ptr), "aclrtFree y failed");
	195	+ CHECK_ACL(aclrtFree(z_ptr), "aclrtFree z failed");
	196	+ }
	197	+
	198	+ (void)session.RemoveGraph(kGraphId);
	199	+ }
	200	+
	201	+ CHECK_ACL(aclrtDestroyStream(stream), "aclrtDestroyStream failed");
	202	+ (void)ge::GEFinalize();
	203	+ return ret_code;
	204	+}

Aexamples/custom_op/tilelang_add_custom_offline/CMakeLists.txt+170-0

@@ -0,0 +1,170 @@
	1	+cmake_minimum_required(VERSION 3.14)
	2	+project(tilelang_add_custom_offline LANGUAGES CXX)
	3	+
	4	+option(TILELANG_OFFLINE_BUILD_CUSTOM_OP "Build libcust_opapi.so" ON)
	5	+option(TILELANG_OFFLINE_BUILD_GRAPH_BUILD "Build graph_build" ON)
	6	+option(TILELANG_OFFLINE_BUILD_MODEL_EXEC "Build model_exec" ON)
	7	+
	8	+set(CMAKE_CXX_STANDARD 17)
	9	+set(CMAKE_CXX_STANDARD_REQUIRED ON)
	10	+set(CMAKE_CXX_EXTENSIONS OFF)
	11	+
	12	+if(NOT CMAKE_BUILD_TYPE)
	13	+ set(CMAKE_BUILD_TYPE Release CACHE STRING "Build type" FORCE)
	14	+endif()
	15	+
	16	+set(COMMON_COMPILE_OPTIONS
	17	+ -Wall
	18	+ -Wextra
	19	+ -Wno-unused-parameter
	20	+)
	21	+
	22	+set(PROJECT_OUTPUT_DIR "${CMAKE_SOURCE_DIR}/output")
	23	+file(MAKE_DIRECTORY "${PROJECT_OUTPUT_DIR}")
	24	+
	25	+if(CMAKE_SYSTEM_NAME STREQUAL "Windows")
	26	+ set(OPP_OS_TYPE "windows")
	27	+else()
	28	+ set(OPP_OS_TYPE "linux")
	29	+endif()
	30	+
	31	+string(TOLOWER "${CMAKE_SYSTEM_PROCESSOR}" CMAKE_SYSTEM_PROCESSOR_LOWER)
	32	+if(CMAKE_SYSTEM_PROCESSOR_LOWER MATCHES "^(aarch64\|arm64)$")
	33	+ set(OPP_CPU_TYPE "aarch64")
	34	+elseif(CMAKE_SYSTEM_PROCESSOR_LOWER MATCHES "^(x86_64\|amd64)$")
	35	+ set(OPP_CPU_TYPE "x86_64")
	36	+else()
	37	+ set(OPP_CPU_TYPE "${CMAKE_SYSTEM_PROCESSOR_LOWER}")
	38	+endif()
	39	+
	40	+set(CUSTOM_OP_OUTPUT_DIR "${PROJECT_OUTPUT_DIR}/op_graph/lib/${OPP_OS_TYPE}/${OPP_CPU_TYPE}")
	41	+file(MAKE_DIRECTORY "${CUSTOM_OP_OUTPUT_DIR}")
	42	+set(CUSTOM_OP_INCLUDE_DIR "${PROJECT_OUTPUT_DIR}/op_graph/include")
	43	+file(MAKE_DIRECTORY "${CUSTOM_OP_INCLUDE_DIR}")
	44	+
	45	+configure_file("${CMAKE_SOURCE_DIR}/ge/add_custom.h" "${CUSTOM_OP_INCLUDE_DIR}/add_custom.h" COPYONLY)
	46	+
	47	+set(KERNEL_SOURCE_FILE "${CMAKE_SOURCE_DIR}/add_custom_kernel/add_custom_kernel.py")
	48	+
	49	+set(ASCEND_HOME_PATH_OVERRIDE "" CACHE PATH "Optional ASCEND_HOME_PATH override")
	50	+if(ASCEND_HOME_PATH_OVERRIDE)
	51	+ set(ASCEND_HOME_PATH "${ASCEND_HOME_PATH_OVERRIDE}")
	52	+else()
	53	+ set(ASCEND_HOME_PATH "$ENV{ASCEND_HOME_PATH}")
	54	+endif()
	55	+
	56	+if(ASCEND_HOME_PATH)
	57	+ message(STATUS "ASCEND_HOME_PATH: ${ASCEND_HOME_PATH}")
	58	+else()
	59	+ message(WARNING "ASCEND_HOME_PATH is empty. Configure succeeds, but compilation requires a valid CANN toolkit path.")
	60	+endif()
	61	+
	62	+if(TILELANG_OFFLINE_BUILD_CUSTOM_OP)
	63	+ configure_file("${KERNEL_SOURCE_FILE}" "${CUSTOM_OP_OUTPUT_DIR}/add_custom_kernel.py" COPYONLY)
	64	+
	65	+ add_library(cust_opapi SHARED
	66	+ ge/custom_op.cpp
	67	+ )
	68	+ target_compile_options(cust_opapi PRIVATE
	69	+ ${COMMON_COMPILE_OPTIONS}
	70	+ )
	71	+ target_compile_definitions(cust_opapi PRIVATE
	72	+ _GLIBCXX_USE_CXX11_ABI=0
	73	+ )
	74	+
	75	+ if(ASCEND_HOME_PATH)
	76	+ target_include_directories(cust_opapi PRIVATE
	77	+ "${CMAKE_SOURCE_DIR}/ge"
	78	+ "${ASCEND_HOME_PATH}/include"
	79	+ "${ASCEND_HOME_PATH}/include/graph"
	80	+ "${ASCEND_HOME_PATH}/include/register"
	81	+ "${ASCEND_HOME_PATH}/include/external"
	82	+ )
	83	+ target_link_directories(cust_opapi PRIVATE
	84	+ "${ASCEND_HOME_PATH}/lib64"
	85	+ )
	86	+ target_link_libraries(cust_opapi PRIVATE
	87	+ -Wl,--no-as-needed
	88	+ ascendcl
	89	+ lowering
	90	+ register
	91	+ gert
	92	+ custom_op_registry_static
	93	+ dl
	94	+ -Wl,--as-needed
	95	+ )
	96	+ endif()
	97	+
	98	+ set_target_properties(cust_opapi PROPERTIES
	99	+ OUTPUT_NAME "cust_opapi"
	100	+ LIBRARY_OUTPUT_DIRECTORY "${CUSTOM_OP_OUTPUT_DIR}"
	101	+ RUNTIME_OUTPUT_DIRECTORY "${CUSTOM_OP_OUTPUT_DIR}"
	102	+ ARCHIVE_OUTPUT_DIRECTORY "${CUSTOM_OP_OUTPUT_DIR}"
	103	+ )
	104	+
	105	+ install(FILES "${CUSTOM_OP_OUTPUT_DIR}/add_custom_kernel.py"
	106	+ DESTINATION "${CUSTOM_OP_OUTPUT_DIR}"
	107	+ )
	108	+ install(FILES "${CUSTOM_OP_INCLUDE_DIR}/add_custom.h"
	109	+ DESTINATION "${CUSTOM_OP_INCLUDE_DIR}"
	110	+ )
	111	+endif()
	112	+
	113	+if(TILELANG_OFFLINE_BUILD_GRAPH_BUILD)
	114	+ add_executable(tilelang_offline_graph_build
	115	+ graph_build/main.cc
	116	+ )
	117	+ target_compile_options(tilelang_offline_graph_build PRIVATE ${COMMON_COMPILE_OPTIONS})
	118	+ target_compile_definitions(tilelang_offline_graph_build PRIVATE
	119	+ _GLIBCXX_USE_CXX11_ABI=0
	120	+ )
	121	+
	122	+ if(ASCEND_HOME_PATH)
	123	+ target_include_directories(tilelang_offline_graph_build PRIVATE
	124	+ "${CUSTOM_OP_INCLUDE_DIR}"
	125	+ "${CMAKE_SOURCE_DIR}/ge"
	126	+ "${ASCEND_HOME_PATH}/include"
	127	+ "${ASCEND_HOME_PATH}/include/graph"
	128	+ "${ASCEND_HOME_PATH}/include/ge"
	129	+ "${ASCEND_HOME_PATH}/opp/built-in/op_proto/inc"
	130	+ )
	131	+ target_link_directories(tilelang_offline_graph_build PRIVATE "${ASCEND_HOME_PATH}/lib64")
	132	+ target_link_libraries(tilelang_offline_graph_build PRIVATE
	133	+ graph
	134	+ ge_runner
	135	+ ge_compiler
	136	+ ascendcl
	137	+ graph_base
	138	+ )
	139	+ endif()
	140	+
	141	+ set_target_properties(tilelang_offline_graph_build PROPERTIES
	142	+ RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}"
	143	+ )
	144	+endif()
	145	+
	146	+if(TILELANG_OFFLINE_BUILD_MODEL_EXEC)
	147	+ add_executable(tilelang_offline_model_exec
	148	+ model_exec/main.cc
	149	+ )
	150	+ target_compile_options(tilelang_offline_model_exec PRIVATE ${COMMON_COMPILE_OPTIONS})
	151	+ target_compile_definitions(tilelang_offline_model_exec PRIVATE
	152	+ _GLIBCXX_USE_CXX11_ABI=0
	153	+ )
	154	+
	155	+ if(ASCEND_HOME_PATH)
	156	+ target_include_directories(tilelang_offline_model_exec PRIVATE
	157	+ "${ASCEND_HOME_PATH}/include"
	158	+ "${ASCEND_HOME_PATH}/include/acl"
	159	+ )
	160	+ target_link_directories(tilelang_offline_model_exec PRIVATE "${ASCEND_HOME_PATH}/lib64")
	161	+ target_link_libraries(tilelang_offline_model_exec PRIVATE
	162	+ ascendcl
	163	+ acl_rt
	164	+ )
	165	+ endif()
	166	+
	167	+ set_target_properties(tilelang_offline_model_exec PROPERTIES
	168	+ RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}"
	169	+ )
	170	+endif()

Aexamples/custom_op/tilelang_add_custom_offline/README.md+212-0

@@ -0,0 +1,212 @@
	1	+# TileLang Add Custom Operator 离线 OM 模型下沉样例
	2	+
	3	+## 样例概述
	4	+
	5	+- 构图入口: GE 原生 (`Graph::SaveToFile` 生成 AIR + ATC 编译 OM)
	6	+- 算子编程语言: TileLang
	7	+- 编译方式: ATC 编译阶段通过 `CompilableOp::Compile` 回调 subprocess 调用 TileLang Python 编译器，在线编译 kernel 源码为 `.so`，再通过 `PortableOp::Serialize` 将 `.so` 字节序列化到 OM 模型
	8	+- 核心链路: `Graph → AIR → ATC 编译(Compile + Serialize) → OM → ACL 加载(Deserialize + Execute)`
	9	+- 场景: 场景 C — 离线 OM 模型下沉（`CompilableOp` + `PortableOp` + `EagerExecuteOp` + `ShapeInferOp`）
	10	+
	11	+本样例以 element-wise Add 算子为例，展示如何通过 `PortableOp` 接口将 TileLang 编译产物序列化到 OM 模型文件，实现离线部署链路。与 [tilelang_add_custom_online](../tilelang_add_custom_online/README.md)（在线编译场景 B）形成对比。
	12	+
	13	+## 与在线编译样例的区别
	14	+
	15	+\| 维度 \| 在线编译 (`tilelang_add_custom_online`) \| 离线 OM 下沉 (本样例) \|
	16	+\|------\|---------------------------------------\|---------------------\|
	17	+\| 接口组合 \| `CompilableOp` + `EagerExecuteOp` + `ShapeInferOp` \| + `PortableOp` \|
	18	+\| 模型格式 \| 无 OM，直接 `Session::ExecuteGraphWithStreamAsync` \| OM 模型文件 \|
	19	+\| 编译产物生命周期 \| 进程内缓存，进程退出即丢失 \| 序列化到 OM 文件，跨进程持久化 \|
	20	+\| 执行方式 \| GE Session 在线执行 \| ACL `aclmdlLoadFromFile` + `aclmdlExecute` \|
	21	+\| 部署方式 \| 需要运行环境具备 Python + TileLang \| OM 文件自包含，部署环境无需 Python + TileLang \|
	22	+
	23	+## 目录结构
	24	+
	25	+```text
	26	+tilelang_add_custom_offline/
	27	+├── README.md
	28	+├── README_en.md
	29	+├── CMakeLists.txt # 构建 libcust_opapi.so + graph_build + model_exec
	30	+├── run.sh # 一键编译运行
	31	+├── add_custom_kernel/
	32	+│ └── add_custom_kernel.py # TileLang kernel 源码（接受 N 和输出路径参数）
	33	+├── ge/
	34	+│ ├── add_custom.h # REG_OP 算子 proto 定义
	35	+│ └── custom_op.cpp # CompilableOp + PortableOp + EagerExecuteOp + ShapeInferOp 实现
	36	+├── graph_build/
	37	+│ └── main.cc # 构图 + Graph::SaveToFile 生成 AIR（供 ATC 编译）
	38	+└── model_exec/
	39	+ └── main.cc # ACL 加载 OM + 执行 + 精度校验（触发 Deserialize + Execute）
	40	+```
	41	+
	42	+## 核心流程
	43	+
	44	+```text
	45	+=== graph_build 阶段 (Graph::SaveToFile → ATC) ===
	46	+
	47	+graph_build 生成 AIR 文件 → ATC 加载 AIR 编译 OM
	48	+
	49	+GE 回调 Compile(ctx)
	50	+ ├─ 读取输入元素数量 → 构建 binary key
	51	+ ├─ exec("python3 add_custom_kernel.py <N> <output.so>")（同机有卡编译）
	52	+ ├─ TileLang 编译器编译 kernel 源码 → 产出 .so（host-wrapper）
	53	+ ├─ 读取 .so 文件字节 → so_data
	54	+ ├─ mkstemps 临时文件读取后立即 unlink
	55	+ └─ dlopen .so + dlsym("call") → 缓存函数指针
	56	+
	57	+GE 回调 Serialize(buffer)
	58	+ ├─ 小端格式: [magic][version][count]
	59	+ │ [key_len][key][so_size][so_data] ...
	60	+ └─ 将 kernel_entries_ 中所有 .so 字节写入 buffer → 嵌入 OM
	61	+
	62	+aclgrphSaveModel → 保存 OM 文件
	63	+
	64	+=== model_exec 阶段 (aclmdlLoadFromFile) ===
	65	+
	66	+ACL 加载 OM → GE 回调 Deserialize(buffer)
	67	+ ├─ 校验 magic/version/count，检查边界和重复 key
	68	+ ├─ 逐条恢复 kernel entry，使用 memfd_create 从内存加载 .so（不落盘）
	69	+ ├─ dlopen memfd + dlsym("call") → 缓存函数指针
	70	+ ├─ 检查尾部无脏数据
	71	+ └─ 全部成功后原子替换 kernel_entries_（事务式）
	72	+
	73	+aclmdlExecute → GE 回调 Execute(ctx)
	74	+ ├─ 从 kernel_entries_ 获取 call 函数指针
	75	+ ├─ 分配输出 Tensor
	76	+ └─ call(x_ptr, y_ptr, z_ptr, stream) → NPU 执行
	77	+```
	78	+
	79	+## 前置依赖
	80	+
	81	+### CANN
	82	+
	83	+- 已正确安装并配置 CANN 环境（`source ${ASCEND_HOME_PATH}/set_env.sh`）
	84	+
	85	+### TileLang-Ascend
	86	+
	87	+需安装 TileLang 主包和 TileLang-Ascend 后端：
	88	+
	89	+```bash
	90	+pip install tilelang
	91	+# TileLang-Ascend 后端：从 https://github.com/tile-ai/tilelang-ascend 安装
	92	+```
	93	+
	94	+> 注意：TileLang-Ascend 仅在 graph_build 阶段（编译 OM）需要，model_exec 阶段（加载执行 OM）不需要。
	95	+
	96	+### 环境变量
	97	+
	98	+\| 变量 \| 必需 \| 说明 \|
	99	+\|------\|------\|------\|
	100	+\| `ASCEND_HOME_PATH` \| 是 \| CANN toolkit 路径 \|
	101	+\| `TILELANG_ASCEND_HOME` \| 否 \| TileLang-Ascend 源码安装路径（pip 安装则无需设置） \|
	102	+\| `ASCEND_CUSTOM_OPP_PATH` \| 自动 \| 由 `run.sh` 自动设置 \|
	103	+
	104	+## 快速运行
	105	+
	106	+```bash
	107	+source ${ASCEND_HOME_PATH}/set_env.sh
	108	+bash run.sh
	109	+```
	110	+
	111	+`run.sh` 依次执行 4 个步骤：
	112	+
	113	+1. 构建 `libcust_opapi.so`、`graph_build` 和 `model_exec`，安装 `.py` 源码到 OPP 包
	114	+2. 运行 `graph_build`（`Graph::SaveToFile` 生成 AIR 文件）
	115	+3. 运行 `atc`（编译 AIR → OM，触发 `Compile` + `Serialize`）
	116	+4. 运行 `model_exec`（`aclmdlLoadFromFile` 触发 `Deserialize`，`aclmdlExecute` 触发 `Execute`）
	117	+
	118	+成功时终端输出：
	119	+
	120	+```text
	121	+[INFO] Step 1/4: build custom op library, graph_build and model_exec
	122	+...
	123	+[INFO] Step 2/4: generate AIR file (graph definition)
	124	+Saving AIR file (for ATC offline compilation)...
	125	+AIR file saved to: .../tilelang_add_offline.air
	126	+[INFO] Step 3/4: compile AIR to OM via ATC (triggers Compile + Serialize)
	127	+ATC compiling ...
	128	+Compiling TileLang kernel: python3 ".../add_custom_kernel.py" 4096 "..." 2>&1
	129	+TileLang kernel compiled and loaded, key=4096, so_size=...
	130	+Serialized 1 kernel(s), total buffer size=...
	131	+[INFO] OM model generated: ... bytes
	132	+[INFO] Step 4/4: execute OM model (triggers Deserialize + Execute)
	133	+Loading OM model (triggers Deserialize): .../tilelang_add_offline.om
	134	+Deserialized 1 kernel(s)
	135	+Executing model (triggers Execute)...
	136	+Precision check passed, max_error=0
	137	+[INFO] Sample pipeline finished.
	138	+```
	139	+
	140	+## 算子规格
	141	+
	142	+\| 项目 \| 内容 \|
	143	+\|------\|------\|
	144	+\| 算子类型 \| `AddCustomOffline` \|
	145	+\| 输入 \| `x` (float32), `y` (float32) \|
	146	+\| 输出 \| `z` (float32) \|
	147	+\| 输入 shape \| `[4096]` (固定) \|
	148	+\| 输出 shape \| `[4096]` \|
	149	+\| 格式 \| ND \|
	150	+\| kernel 名称 \| `main_kernel`（由 `call` 封装） \|
	151	+\| BLOCK_SIZE \| 1024 \|
	152	+
	153	+## 序列化格式
	154	+
	155	+`PortableOp::Serialize` 使用自定义二进制格式将 TileLang `.so` 编译产物嵌入 OM：
	156	+
	157	+```text
	158	+偏移长度字段说明
	159	+0 4 magic 固定 0x4F504B4E（自定义格式标识，小端）
	160	+4 4 version 固定 1（小端）
	161	+8 4 count kernel 条目数（小端）
	162	+12 --- entries 重复 count 次:
	163	+ 4 key_len key 字节长度（小端）
	164	+ N key 元素数量字符串（如 "4096"）
	165	+ 4 so_size .so 文件字节长度（小端）
	166	+ M so_data .so 完整二进制内容
	167	+```
	168	+
	169	+`Deserialize` 读取该格式，对每个 `.so` 使用 `memfd_create` 从内存加载（不落盘），并检查重复 key、尾部脏数据等完整性约束。
	170	+
	171	+## 关键文件说明
	172	+
	173	+### `ge/custom_op.cpp`
	174	+
	175	+GE 交付件，实现 `CompilableOp` + `PortableOp` + `EagerExecuteOp` + `ShapeInferOp`：
	176	+
	177	+- Compile: subprocess 调用 Python 编译 TileLang → 读取 `.so` 字节 → `dlopen` → 缓存
	178	+- Serialize: 将 `kernel_entries_` 中的 `.so` 字节序列化为二进制 buffer
	179	+- Deserialize: 从 buffer 恢复 `.so` 字节 → 写临时文件 → `dlopen` → 缓存
	180	+- Execute: 使用缓存的函数指针调用 `call(x, y, z, stream)`
	181	+- 使用 `std::mutex` 保证线程安全
	182	+
	183	+### `graph_build/main.cc`
	184	+
	185	+使用 `Graph::SaveToFile` 生成 AIR 文件，供 ATC 离线编译：
	186	+
	187	+1. `GEInitialize` + 构建计算图
	188	+2. `graph->SaveToFile(air_path)` — 生成 AIR 文件
	189	+
	190	+ATC 编译 AIR → OM 时会自动触发 `Compile` + `Serialize`。
	191	+
	192	+### `model_exec/main.cc`
	193	+
	194	+使用 ACL API 加载和执行 OM 模型：
	195	+
	196	+1. `aclInit` + `aclrtSetDevice`
	197	+2. `aclmdlLoadFromFile(om_path)` — 触发 Deserialize
	198	+3. `aclmdlGetDesc` 获取模型描述
	199	+4. 分配 device 内存，H2D 拷贝输入
	200	+5. `aclmdlExecute` — 触发 Execute
	201	+6. D2H 拷贝输出，精度校验
	202	+
	203	+## 注意事项
	204	+
	205	+- 同机有卡编译限定：TileLang-Ascend 当前通过 `torch.npu.get_device_name()` 做运行时平台检测，不支持离线指定目标架构。因此 `Compile` 回调中调用 Python 编译器时未传递 `--soc_version`，编译产物绑定编译机 NPU。本样例仅适用于"编译机与目标机为同一 NPU"的场景，不能用于跨平台 ATC 离线编译。如需跨平台编译，需等待 TileLang 支持离线目标指定后更新。
	206	+- `graph_build` 阶段需要运行环境具备 Python + TileLang-Ascend；`model_exec` 阶段不需要（OM 自包含 `.so` 编译产物）。
	207	+- `ge.graphRunMode=1` 确保走在线执行链路（PRIORITY_GRAPH 模式）。
	208	+- 序列化格式为自定义格式，GE 只透传不解析，格式完全由算子控制。所有 `uint32_t` 字段使用小端格式。
	209	+- 当前样例仅支持 float32，如需支持更多数据类型需调整 `REG_OP` 的 `DATATYPE` 约束和 TileLang kernel 的 dtype 参数。
	210	+- OM 编译依赖 ATC 工具：`graph_build` 生成 AIR 文件，ATC 负责编译 AIR → OM（触发 `Compile` + `Serialize`）。`SOC_VERSION` 环境变量可覆盖默认的 `Ascend910_9362`。
	211	+- 编译产物 `.so` 在 `Compile` 阶段使用 `mkstemps` 生成唯一临时文件并在读取后立即 `unlink`；`Deserialize` 阶段使用 `memfd_create` 从内存加载，不落盘。
	212	+- 若当前环境 ATC 不可用（版本不匹配等），可参考 `compilable_add_custom` 样例在 ATC 可用的环境中编译 OM。

Aexamples/custom_op/tilelang_add_custom_offline/README_en.md+47-0

@@ -0,0 +1,47 @@
	1	+# TileLang Add Custom Operator Offline OM Model Sinking Sample
	2	+
	3	+## Sample Overview
	4	+
	5	+- Graph construction entry: GE native (`Graph::SaveToFile` generates AIR, then ATC compiles OM)
	6	+- Operator programming language: TileLang
	7	+- Compilation method: ATC compile phase invokes TileLang Python compiler via `CompilableOp::Compile` callback (subprocess), then `PortableOp::Serialize` embeds `.so` bytes into OM model
	8	+- Core pipeline: `Graph → AIR → ATC compile (Compile + Serialize) → OM → ACL load (Deserialize + Execute)`
	9	+- Scenario: Scenario C — offline OM model sinking (`CompilableOp` + `PortableOp` + `EagerExecuteOp` + `ShapeInferOp`)
	10	+
	11	+This sample demonstrates how to serialize TileLang compilation products into an OM model file via the `PortableOp` interface, enabling offline deployment. Contrast with [tilelang_add_custom_online](../tilelang_add_custom_online/README_en.md) (online compilation, Scenario B).
	12	+
	13	+## Differences from Online Compilation Sample
	14	+
	15	+\| Dimension \| Online (`tilelang_add_custom_online`) \| Offline OM (this sample) \|
	16	+\|-----------\|---------------------------------------\|--------------------------\|
	17	+\| Interface combo \| `CompilableOp` + `EagerExecuteOp` + `ShapeInferOp` \| + `PortableOp` \|
	18	+\| Model format \| No OM, direct `Session::ExecuteGraphWithStreamAsync` \| OM model file \|
	19	+\| Compilation product lifecycle \| In-process cache, lost on process exit \| Serialized to OM file, persists across processes \|
	20	+\| Execution \| GE Session online execution \| ACL `aclmdlLoadFromFile` + `aclmdlExecute` \|
	21	+\| Deployment \| Requires Python + TileLang in runtime \| OM file is self-contained, no Python + TileLang needed at deployment \|
	22	+
	23	+## Quick Start
	24	+
	25	+```bash
	26	+source ${ASCEND_HOME_PATH}/set_env.sh
	27	+bash run.sh
	28	+```
	29	+
	30	+## Operator Specification
	31	+
	32	+\| Item \| Value \|
	33	+\|------\|-------\|
	34	+\| Op type \| `AddCustomOffline` \|
	35	+\| Inputs \| `x` (float32), `y` (float32) \|
	36	+\| Output \| `z` (float32) \|
	37	+\| Input shape \| `[4096]` (fixed) \|
	38	+\| Format \| ND \|
	39	+\| Kernel name \| `main_kernel` (wrapped by `call`) \|
	40	+\| BLOCK_SIZE \| 1024 \|
	41	+
	42	+## Notes
	43	+
	44	+- Same-machine NPU compilation required: TileLang-Ascend uses `torch.npu.get_device_name()` for runtime platform detection and does not support specifying target architecture offline. This sample only works when the compilation machine has the same NPU as the target. Cross-platform ATC offline compilation is not supported until TileLang adds offline target specification.
	45	+- `graph_build` phase requires Python + TileLang-Ascend; `model_exec` phase does not (OM is self-contained).
	46	+- Serialization format is custom (little-endian); GE only transparently passes the buffer.
	47	+- `Deserialize` uses `memfd_create` to load `.so` from memory (no disk files), with boundary checks, duplicate key detection, trailing data validation, and transactional rollback on failure.

Aexamples/custom_op/tilelang_add_custom_offline/add_custom_kernel/add_custom_kernel.py+105-0

@@ -0,0 +1,105 @@
	1	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	2	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	3	+# CANN Open Software License Agreement Version 2.0 (the "License").
	4	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	5	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	6	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	7	+# See LICENSE in the root of the software repository for the full text of the License.
	8	+
	9	+"""
	10	+TileLang Add Kernel + 编译产出 .so
	11	+
	12	+本文件用 TileLang 实现 element-wise Add kernel，并编译产出 Ascend .so 交付件。
	13	+供 GE 自定义算子 (CompilableOp::Compile) 在 GE/ATC 编译阶段通过 subprocess 调用本脚本完成在线编译。
	14	+
	15	+用法:
	16	+ python3 add_custom_kernel.py [N] [output_path]
	17	+
	18	+参数:
	19	+ N - 元素总数（默认 4096，需为 BLOCK_SIZE 的整数倍）
	20	+ output_path - 产出 .so 的路径（默认脚本目录下 add_kernel.so）
	21	+
	22	+TileLang-Ascend 编译后的 .so 导出函数签名为:
	23	+ extern "C" void call(uint8_t* A_handle, uint8_t* B_handle, uint8_t* C_handle, aclrtStream stream)
	24	+内部封装了 main_kernel<<<>>> 的 launch 逻辑。
	25	+"""
	26	+
	27	+import os
	28	+import shutil
	29	+import sys
	30	+
	31	+import tilelang
	32	+import tilelang.language as T
	33	+
	34	+BLOCK_SIZE = 1024
	35	+
	36	+
	37	+@tilelang.jit(out_idx=[-1])
	38	+def vec_add(n, block_size, dtype="float"):
	39	+ m_num = n // block_size
	40	+ vec_num = 2
	41	+
	42	+ @T.prim_func
	43	+ def main(
	44	+ a: T.Tensor((n,), dtype),
	45	+ b: T.Tensor((n,), dtype),
	46	+ c: T.Tensor((n,), dtype),
	47	+ ):
	48	+ with T.Kernel(m_num, is_npu=True) as (cid, vid):
	49	+ a_ub = T.alloc_ub((block_size // vec_num,), dtype)
	50	+ b_ub = T.alloc_ub((block_size // vec_num,), dtype)
	51	+ c_ub = T.alloc_ub((block_size // vec_num,), dtype)
	52	+ with T.Scope("V"):
	53	+ T.copy(a[cid * block_size + vid * block_size // vec_num], a_ub)
	54	+ T.copy(b[cid * block_size + vid * block_size // vec_num], b_ub)
	55	+
	56	+ T.barrier_all()
	57	+ T.tile.add(c_ub, a_ub, b_ub)
	58	+ T.barrier_all()
	59	+
	60	+ T.copy(c_ub, c[cid * block_size + vid * block_size // vec_num])
	61	+
	62	+ return main
	63	+
	64	+
	65	+def main():
	66	+ n = int(sys.argv[1]) if len(sys.argv) > 1 else 4096
	67	+ output_path = sys.argv[2] if len(sys.argv) > 2 else None
	68	+
	69	+ if n % BLOCK_SIZE != 0:
	70	+ print(
	71	+ f"Error: N={n} must be a multiple of BLOCK_SIZE={BLOCK_SIZE}",
	72	+ file=sys.stderr,
	73	+ )
	74	+ sys.exit(1)
	75	+
	76	+ func = vec_add(n, BLOCK_SIZE)
	77	+
	78	+ adapter = func.adapter
	79	+ so_path = getattr(getattr(adapter, "lib", None), "_name", None)
	80	+ if so_path is None:
	81	+ print(
	82	+ "Error: cannot locate compiled .so path from tilelang adapter",
	83	+ file=sys.stderr,
	84	+ )
	85	+ sys.exit(1)
	86	+
	87	+ if output_path is None:
	88	+ output_path = os.path.join(
	89	+ os.path.dirname(os.path.abspath(__file__)), "add_kernel.so"
	90	+ )
	91	+
	92	+ output_dir = os.path.dirname(os.path.abspath(output_path))
	93	+ os.makedirs(output_dir, exist_ok=True)
	94	+ shutil.copy2(so_path, output_path)
	95	+
	96	+ print(f"Kernel .so saved to: {output_path}")
	97	+ print(f"N={n}, BLOCK_SIZE={BLOCK_SIZE}")
	98	+ print(f"File size: {os.path.getsize(output_path)} bytes")
	99	+ print(
	100	+ "Export function: call(uint8_t* A, uint8_t* B, uint8_t* C, aclrtStream stream)"
	101	+ )
	102	+
	103	+
	104	+if __name__ == "__main__":
	105	+ main()

Aexamples/custom_op/tilelang_add_custom_offline/ge/add_custom.h+24-0

@@ -0,0 +1,24 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#ifndef EXAMPLES_CUSTOM_OP_TILELANG_ADD_CUSTOM_OFFLINE_GE_ADD_CUSTOM_H_
	12	+#define EXAMPLES_CUSTOM_OP_TILELANG_ADD_CUSTOM_OFFLINE_GE_ADD_CUSTOM_H_
	13	+
	14	+#include "graph/operator_reg.h"
	15	+
	16	+namespace ge {
	17	+REG_OP(AddCustomOffline)
	18	+ .INPUT(x, TensorType({DT_FLOAT}))
	19	+ .INPUT(y, TensorType({DT_FLOAT}))
	20	+ .OUTPUT(z, TensorType({DT_FLOAT}))
	21	+ .OP_END_FACTORY_REG(AddCustomOffline);
	22	+} // namespace ge
	23	+
	24	+#endif // EXAMPLES_CUSTOM_OP_TILELANG_ADD_CUSTOM_OFFLINE_GE_ADD_CUSTOM_H_

Aexamples/custom_op/tilelang_add_custom_offline/ge/custom_op.cpp+455-0

@@ -0,0 +1,455 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#include <cstdint>
	12	+#include <cstdio>
	13	+#include <cstdlib>
	14	+#include <cstring>
	15	+#include <dlfcn.h>
	16	+#include <fcntl.h>
	17	+#include <fstream>
	18	+#include <iostream>
	19	+#include <map>
	20	+#include <mutex>
	21	+#include <set>
	22	+#include <sstream>
	23	+#include <string>
	24	+#include <sys/mman.h>
	25	+#include <sys/stat.h>
	26	+#include <sys/wait.h>
	27	+#include <unistd.h>
	28	+#include "graph/custom_op.h"
	29	+#include "acl/acl_rt.h"
	30	+
	31	+using namespace ge;
	32	+
	33	+namespace {
	34	+constexpr const char *kKernelSourceFile = "add_custom_kernel.py";
	35	+constexpr const char *kCallFuncName = "call";
	36	+constexpr uint32_t kSerializeMagic = 0x4F504B4EU; // "OPKN" in ASCII, custom format identifier
	37	+constexpr uint32_t kSerializeVersion = 1U;
	38	+constexpr size_t kMaxSoSize = 100U * 1024U * 1024U; // 100 MB
	39	+constexpr size_t kMaxKeyLen = 256U;
	40	+constexpr size_t kMaxEntryCount = 64U;
	41	+
	42	+using CallFunc = void ()(void x_ptr, void y_ptr, void z_ptr, void *stream);
	43	+
	44	+struct KernelEntry {
	45	+ std::vector<uint8_t> so_data;
	46	+ void *handle = nullptr;
	47	+ CallFunc call_func = nullptr;
	48	+};
	49	+
	50	+using KernelEntryMap = std::map<std::string, KernelEntry>;
	51	+
	52	+std::string GetCurrentLibraryDir() {
	53	+ Dl_info info{};
	54	+ if ((dladdr(reinterpret_cast<void *>(&GetCurrentLibraryDir), &info) == 0) \|\| (info.dli_fname == nullptr)) {
	55	+ return {};
	56	+ }
	57	+ const std::string library_path = info.dli_fname;
	58	+ const auto pos = library_path.find_last_of('/');
	59	+ if (pos == std::string::npos) {
	60	+ return ".";
	61	+ }
	62	+ if (pos == 0U) {
	63	+ return "/";
	64	+ }
	65	+ return library_path.substr(0U, pos);
	66	+}
	67	+
	68	+std::string GetKernelSourcePath() {
	69	+ const auto library_dir = GetCurrentLibraryDir();
	70	+ if (library_dir.empty()) {
	71	+ return {};
	72	+ }
	73	+ if (library_dir == "/") {
	74	+ return library_dir + kKernelSourceFile;
	75	+ }
	76	+ return library_dir + "/" + kKernelSourceFile;
	77	+}
	78	+
	79	+std::string BuildBinaryKey(int64_t shape_size) {
	80	+ return std::to_string(shape_size);
	81	+}
	82	+
	83	+bool ReadFileToVector(const std::string &path, std::vector<uint8_t> &data) {
	84	+ std::ifstream file(path, std::ios::binary);
	85	+ if (!file) {
	86	+ return false;
	87	+ }
	88	+ std::ostringstream buffer;
	89	+ buffer << file.rdbuf();
	90	+ const auto &str = buffer.str();
	91	+ data.assign(str.begin(), str.end());
	92	+ return true;
	93	+}
	94	+
	95	+void WriteU32LE(std::vector<uint8_t> &buf, uint32_t val) {
	96	+ buf.push_back(static_cast<uint8_t>(val & 0xFFU));
	97	+ buf.push_back(static_cast<uint8_t>((val >> 8U) & 0xFFU));
	98	+ buf.push_back(static_cast<uint8_t>((val >> 16U) & 0xFFU));
	99	+ buf.push_back(static_cast<uint8_t>((val >> 24U) & 0xFFU));
	100	+}
	101	+
	102	+bool ReadU32LE(const std::vector<uint8_t> &buf, size_t &offset, uint32_t &val) {
	103	+ if (offset + sizeof(uint32_t) > buf.size()) {
	104	+ return false;
	105	+ }
	106	+ val = static_cast<uint32_t>(buf[offset]) \| (static_cast<uint32_t>(buf[offset + 1U]) << 8U) \|
	107	+ (static_cast<uint32_t>(buf[offset + 2U]) << 16U) \| (static_cast<uint32_t>(buf[offset + 3U]) << 24U);
	108	+ offset += sizeof(uint32_t);
	109	+ return true;
	110	+}
	111	+
	112	+bool ReadBytesSafe(const std::vector<uint8_t> &buf, size_t &offset, size_t len, std::vector<uint8_t> &out) {
	113	+ if (len > buf.size() \|\| offset > buf.size() - len) {
	114	+ return false;
	115	+ }
	116	+ out.assign(buf.data() + offset, buf.data() + offset + len);
	117	+ offset += len;
	118	+ return true;
	119	+}
	120	+
	121	+bool WriteVectorToMemfd(const std::vector<uint8_t> &data, std::string &memfd_path) {
	122	+ int fd = memfd_create("tilelang_kernel", MFD_CLOEXEC);
	123	+ if (fd < 0) {
	124	+ return false;
	125	+ }
	126	+ ssize_t written = 0;
	127	+ const size_t total = data.size();
	128	+ while (static_cast<size_t>(written) < total) {
	129	+ ssize_t ret = write(fd, data.data() + written, total - static_cast<size_t>(written));
	130	+ if (ret < 0) {
	131	+ (void)close(fd);
	132	+ return false;
	133	+ }
	134	+ written += ret;
	135	+ }
	136	+ memfd_path = "/proc/self/fd/" + std::to_string(fd);
	137	+ return true;
	138	+}
	139	+
	140	+graphStatus LoadSoFromData(const std::vector<uint8_t> &so_data, void *&handle, CallFunc &call_func) {
	141	+ std::string memfd_path;
	142	+ if (!WriteVectorToMemfd(so_data, memfd_path)) {
	143	+ std::cerr << "Failed to create memfd for kernel .so" << std::endl;
	144	+ return GRAPH_FAILED;
	145	+ }
	146	+
	147	+ handle = dlopen(memfd_path.c_str(), RTLD_NOW);
	148	+ if (handle == nullptr) {
	149	+ std::cerr << "dlopen failed: " << dlerror() << std::endl;
	150	+ return GRAPH_FAILED;
	151	+ }
	152	+
	153	+ dlerror();
	154	+ call_func = reinterpret_cast<CallFunc>(dlsym(handle, kCallFuncName));
	155	+ const char *error = dlerror();
	156	+ if (error != nullptr) {
	157	+ std::cerr << "dlsym '" << kCallFuncName << "' failed: " << error << std::endl;
	158	+ (void)dlclose(handle);
	159	+ handle = nullptr;
	160	+ return GRAPH_FAILED;
	161	+ }
	162	+ return GRAPH_SUCCESS;
	163	+}
	164	+
	165	+int ExecuteSubprocess(const char *const argv[], std::string &output) {
	166	+ std::ostringstream cmd_stream;
	167	+ for (size_t i = 0U; argv[i] != nullptr; ++i) {
	168	+ if (i > 0U) {
	169	+ cmd_stream << " ";
	170	+ }
	171	+ cmd_stream << argv[i];
	172	+ }
	173	+ const std::string cmd = cmd_stream.str() + " 2>&1";
	174	+
	175	+ FILE *pipe = popen(cmd.c_str(), "r");
	176	+ if (pipe == nullptr) {
	177	+ return -1;
	178	+ }
	179	+ char buffer[256];
	180	+ size_t n = 0U;
	181	+ while ((n = fread(buffer, 1U, sizeof(buffer), pipe)) > 0U) {
	182	+ output.append(buffer, n);
	183	+ }
	184	+ int status = pclose(pipe);
	185	+ if (WIFEXITED(status)) {
	186	+ return WEXITSTATUS(status);
	187	+ }
	188	+ return -1;
	189	+}
	190	+} // namespace
	191	+
	192	+class AddCustomOffline : public CompilableOp, public PortableOp, public EagerExecuteOp, public ShapeInferOp {
	193	+ public:
	194	+ ~AddCustomOffline() {
	195	+ for (auto &entry : kernel_entries_) {
	196	+ if (entry.second.handle != nullptr) {
	197	+ (void)dlclose(entry.second.handle);
	198	+ }
	199	+ }
	200	+ }
	201	+
	202	+ graphStatus Compile(gert::OpCompileContext *ctx) override {
	203	+ if (ctx == nullptr) {
	204	+ std::cerr << "Compile context is null" << std::endl;
	205	+ return GRAPH_FAILED;
	206	+ }
	207	+
	208	+ const gert::Tensor *input_x = ctx->GetInputTensor(0);
	209	+ if (input_x == nullptr) {
	210	+ std::cerr << "Compile: GetInputTensor(0) failed" << std::endl;
	211	+ return GRAPH_FAILED;
	212	+ }
	213	+
	214	+ const int64_t n = input_x->GetShapeSize();
	215	+ const std::string key = BuildBinaryKey(n);
	216	+
	217	+ std::lock_guard<std::mutex> guard(mutex_);
	218	+ if (kernel_entries_.find(key) != kernel_entries_.end()) {
	219	+ std::cout << "TileLang kernel already compiled for key=" << key << std::endl;
	220	+ return GRAPH_SUCCESS;
	221	+ }
	222	+
	223	+ const std::string py_path = GetKernelSourcePath();
	224	+ if (py_path.empty()) {
	225	+ std::cerr << "Failed to locate TileLang kernel source: " << kKernelSourceFile << std::endl;
	226	+ return GRAPH_FAILED;
	227	+ }
	228	+
	229	+ char so_tmpl[] = "/tmp/tilelang_offline_XXXXXX.so";
	230	+ int tmp_fd = mkstemps(so_tmpl, 3);
	231	+ if (tmp_fd < 0) {
	232	+ std::cerr << "Failed to create temp file for kernel .so" << std::endl;
	233	+ return GRAPH_FAILED;
	234	+ }
	235	+ std::string so_path(so_tmpl);
	236	+
	237	+ std::ostringstream n_str;
	238	+ n_str << n;
	239	+ const std::string n_arg = n_str.str();
	240	+ const char *const argv[] = {"python3", py_path.c_str(), n_arg.c_str(), so_path.c_str(), nullptr};
	241	+ std::cout << "Compiling TileLang kernel (same-machine NPU required)" << std::endl;
	242	+
	243	+ std::string output;
	244	+ const int status = ExecuteSubprocess(argv, output);
	245	+ if (status != 0) {
	246	+ std::cerr << "TileLang compilation failed (exit=" << status << "):" << std::endl;
	247	+ std::cerr << output << std::endl;
	248	+ (void)close(tmp_fd);
	249	+ (void)unlink(so_path.c_str());
	250	+ return GRAPH_FAILED;
	251	+ }
	252	+ std::cout << output;
	253	+ (void)close(tmp_fd);
	254	+
	255	+ std::vector<uint8_t> so_data;
	256	+ if (!ReadFileToVector(so_path, so_data)) {
	257	+ std::cerr << "Failed to read compiled .so: " << so_path << std::endl;
	258	+ (void)unlink(so_path.c_str());
	259	+ return GRAPH_FAILED;
	260	+ }
	261	+ (void)unlink(so_path.c_str());
	262	+
	263	+ void *handle = nullptr;
	264	+ CallFunc call_func = nullptr;
	265	+ if (LoadSoFromData(so_data, handle, call_func) != GRAPH_SUCCESS) {
	266	+ return GRAPH_FAILED;
	267	+ }
	268	+
	269	+ kernel_entries_[key] = {std::move(so_data), handle, call_func};
	270	+ std::cout << "TileLang kernel compiled and loaded, key=" << key
	271	+ << ", so_size=" << kernel_entries_[key].so_data.size() << std::endl;
	272	+ return GRAPH_SUCCESS;
	273	+ }
	274	+
	275	+ graphStatus Serialize(std::vector<uint8_t> &buffer) override {
	276	+ std::lock_guard<std::mutex> guard(mutex_);
	277	+ WriteU32LE(buffer, kSerializeMagic);
	278	+ WriteU32LE(buffer, kSerializeVersion);
	279	+ WriteU32LE(buffer, static_cast<uint32_t>(kernel_entries_.size()));
	280	+
	281	+ for (const auto &entry : kernel_entries_) {
	282	+ const auto &key = entry.first;
	283	+ const auto &so_data = entry.second.so_data;
	284	+ WriteU32LE(buffer, static_cast<uint32_t>(key.size()));
	285	+ buffer.insert(buffer.end(), reinterpret_cast<const uint8_t *>(key.data()),
	286	+ reinterpret_cast<const uint8_t *>(key.data()) + key.size());
	287	+ WriteU32LE(buffer, static_cast<uint32_t>(so_data.size()));
	288	+ buffer.insert(buffer.end(), so_data.data(), so_data.data() + so_data.size());
	289	+ }
	290	+
	291	+ std::cout << "Serialized " << kernel_entries_.size() << " kernel(s), total buffer size=" << buffer.size()
	292	+ << std::endl;
	293	+ return GRAPH_SUCCESS;
	294	+ }
	295	+
	296	+ graphStatus Deserialize(const std::vector<uint8_t> &buffer) override {
	297	+ std::lock_guard<std::mutex> guard(mutex_);
	298	+ size_t offset = 0U;
	299	+ uint32_t magic = 0U;
	300	+ uint32_t version = 0U;
	301	+ uint32_t count = 0U;
	302	+
	303	+ if (!ReadU32LE(buffer, offset, magic) \|\| magic != kSerializeMagic) {
	304	+ std::cerr << "Deserialize: invalid magic" << std::endl;
	305	+ return GRAPH_FAILED;
	306	+ }
	307	+ if (!ReadU32LE(buffer, offset, version) \|\| version != kSerializeVersion) {
	308	+ std::cerr << "Deserialize: unsupported version=" << version << std::endl;
	309	+ return GRAPH_FAILED;
	310	+ }
	311	+ if (!ReadU32LE(buffer, offset, count)) {
	312	+ std::cerr << "Deserialize: failed to read count" << std::endl;
	313	+ return GRAPH_FAILED;
	314	+ }
	315	+ if (count == 0U) {
	316	+ std::cerr << "Deserialize: empty entry count" << std::endl;
	317	+ return GRAPH_FAILED;
	318	+ }
	319	+ if (count > kMaxEntryCount) {
	320	+ std::cerr << "Deserialize: entry count " << count << " exceeds limit " << kMaxEntryCount << std::endl;
	321	+ return GRAPH_FAILED;
	322	+ }
	323	+
	324	+ KernelEntryMap temp_entries;
	325	+ std::set<std::string> seen_keys;
	326	+ for (uint32_t i = 0U; i < count; ++i) {
	327	+ if (DeserializeOneEntry(buffer, offset, i, seen_keys, temp_entries) != GRAPH_SUCCESS) {
	328	+ ReleaseHandles(temp_entries);
	329	+ return GRAPH_FAILED;
	330	+ }
	331	+ }
	332	+
	333	+ if (offset != buffer.size()) {
	334	+ std::cerr << "Deserialize: trailing data after " << count << " entries, offset=" << offset
	335	+ << ", buffer_size=" << buffer.size() << std::endl;
	336	+ ReleaseHandles(temp_entries);
	337	+ return GRAPH_FAILED;
	338	+ }
	339	+
	340	+ kernel_entries_ = std::move(temp_entries);
	341	+ std::cout << "Deserialized " << kernel_entries_.size() << " kernel(s)" << std::endl;
	342	+ return GRAPH_SUCCESS;
	343	+ }
	344	+
	345	+ private:
	346	+ void ReleaseHandles(KernelEntryMap &entries) {
	347	+ for (auto &e : entries) {
	348	+ if (e.second.handle != nullptr) {
	349	+ (void)dlclose(e.second.handle);
	350	+ }
	351	+ }
	352	+ }
	353	+
	354	+ graphStatus DeserializeOneEntry(const std::vector<uint8_t> &buffer, size_t &offset, uint32_t index,
	355	+ std::set<std::string> &seen_keys, KernelEntryMap &temp_entries) {
	356	+ uint32_t key_len = 0U;
	357	+ if (!ReadU32LE(buffer, offset, key_len)) {
	358	+ std::cerr << "Deserialize: failed to read key_len at entry " << index << std::endl;
	359	+ return GRAPH_FAILED;
	360	+ }
	361	+ if (key_len == 0U \|\| key_len > kMaxKeyLen) {
	362	+ std::cerr << "Deserialize: invalid key_len " << key_len << " at entry " << index << std::endl;
	363	+ return GRAPH_FAILED;
	364	+ }
	365	+ std::vector<uint8_t> key_bytes;
	366	+ if (!ReadBytesSafe(buffer, offset, key_len, key_bytes)) {
	367	+ std::cerr << "Deserialize: failed to read key at entry " << index << std::endl;
	368	+ return GRAPH_FAILED;
	369	+ }
	370	+ std::string key(key_bytes.begin(), key_bytes.end());
	371	+ if (seen_keys.count(key) > 0) {
	372	+ std::cerr << "Deserialize: duplicate key '" << key << "' at entry " << index << std::endl;
	373	+ return GRAPH_FAILED;
	374	+ }
	375	+
	376	+ uint32_t so_size = 0U;
	377	+ if (!ReadU32LE(buffer, offset, so_size)) {
	378	+ std::cerr << "Deserialize: failed to read so_size at entry " << index << std::endl;
	379	+ return GRAPH_FAILED;
	380	+ }
	381	+ if (so_size == 0U \|\| so_size > kMaxSoSize) {
	382	+ std::cerr << "Deserialize: invalid so_size " << so_size << " at entry " << index << std::endl;
	383	+ return GRAPH_FAILED;
	384	+ }
	385	+ std::vector<uint8_t> so_data;
	386	+ if (!ReadBytesSafe(buffer, offset, so_size, so_data)) {
	387	+ std::cerr << "Deserialize: failed to read so_data at entry " << index << std::endl;
	388	+ return GRAPH_FAILED;
	389	+ }
	390	+
	391	+ void *handle = nullptr;
	392	+ CallFunc call_func = nullptr;
	393	+ if (LoadSoFromData(so_data, handle, call_func) != GRAPH_SUCCESS) {
	394	+ return GRAPH_FAILED;
	395	+ }
	396	+
	397	+ seen_keys.insert(key);
	398	+ temp_entries[key] = {std::move(so_data), handle, call_func};
	399	+ std::cout << "Deserialized kernel, key=" << key << ", so_size=" << temp_entries[key].so_data.size() << std::endl;
	400	+ return GRAPH_SUCCESS;
	401	+ }
	402	+
	403	+ graphStatus Execute(gert::EagerOpExecutionContext *ctx) override {
	404	+ const gert::Tensor *input_x = ctx->GetInputTensor(0);
	405	+ const gert::Tensor *input_y = ctx->GetInputTensor(1);
	406	+ if (input_x == nullptr \|\| input_y == nullptr) {
	407	+ std::cerr << "Execute: GetInputTensor failed" << std::endl;
	408	+ return GRAPH_FAILED;
	409	+ }
	410	+
	411	+ const int64_t n = input_x->GetShapeSize();
	412	+ const std::string key = BuildBinaryKey(n);
	413	+
	414	+ auto it = kernel_entries_.find(key);
	415	+ if (it == kernel_entries_.end()) {
	416	+ std::cerr << "Execute: kernel not found for key=" << key << std::endl;
	417	+ return GRAPH_FAILED;
	418	+ }
	419	+
	420	+ gert::Tensor *output_z =
	421	+ ctx->MallocOutputTensor(0, input_x->GetShape(), input_x->GetFormat(), input_x->GetDataType());
	422	+ if (output_z == nullptr) {
	423	+ std::cerr << "MallocOutputTensor failed" << std::endl;
	424	+ return GRAPH_FAILED;
	425	+ }
	426	+
	427	+ void *stream = ctx->GetStream();
	428	+ it->second.call_func(const_cast<void >(input_x->GetAddr()), const_cast<void >(input_y->GetAddr()),
	429	+ output_z->GetAddr(), stream);
	430	+ return GRAPH_SUCCESS;
	431	+ }
	432	+
	433	+ graphStatus InferShape(gert::InferShapeContext *ctx) override {
	434	+ const auto *input_shape = ctx->GetInputShape(0);
	435	+ auto *output_shape = ctx->GetOutputShape(0);
	436	+ if (input_shape == nullptr \|\| output_shape == nullptr) {
	437	+ return GRAPH_FAILED;
	438	+ }
	439	+ output_shape->SetDimNum(input_shape->GetDimNum());
	440	+ for (size_t i = 0; i < input_shape->GetDimNum(); ++i) {
	441	+ output_shape->SetDim(i, input_shape->GetDim(i));
	442	+ }
	443	+ return GRAPH_SUCCESS;
	444	+ }
	445	+
	446	+ graphStatus InferDataType(gert::InferDataTypeContext *ctx) override {
	447	+ return ctx->SetOutputDataType(0, ctx->GetInputDataType(0));
	448	+ }
	449	+
	450	+ private:
	451	+ std::mutex mutex_;
	452	+ KernelEntryMap kernel_entries_;
	453	+};
	454	+
	455	+REG_AUTO_MAPPING_OP(AddCustomOffline);

Aexamples/custom_op/tilelang_add_custom_offline/graph_build/main.cc+90-0

@@ -0,0 +1,90 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#include <iostream>
	12	+#include <map>
	13	+#include <memory>
	14	+#include <string>
	15	+
	16	+#include "acl/acl_rt.h"
	17	+#include "ge/ge_api.h"
	18	+#include "ge/ge_ir_build.h"
	19	+#include "graph.h"
	20	+#include "ops_proto_legacy.h"
	21	+#include "tensor.h"
	22	+#include "types.h"
	23	+#include "add_custom.h"
	24	+
	25	+using ge::Operator;
	26	+
	27	+namespace {
	28	+constexpr int64_t kNumElements = 4096;
	29	+
	30	+std::unique_ptr<ge::Graph> BuildGraph() {
	31	+ ge::TensorDesc input_desc(ge::Shape({kNumElements}), ge::FORMAT_ND, ge::DT_FLOAT);
	32	+
	33	+ auto data_x = ge::op::Data("data_x");
	34	+ data_x.update_input_desc_x(input_desc);
	35	+ data_x.update_output_desc_y(input_desc);
	36	+ auto data_y = ge::op::Data("data_y");
	37	+ data_y.update_input_desc_x(input_desc);
	38	+ data_y.update_output_desc_y(input_desc);
	39	+
	40	+ auto add = ge::op::AddCustomOffline("add").set_input_x(data_x).set_input_y(data_y);
	41	+ add.update_output_desc_z(input_desc);
	42	+
	43	+ std::vector<Operator> inputs = {data_x, data_y};
	44	+ std::vector<Operator> outputs = {add};
	45	+
	46	+ auto graph = std::make_unique<ge::Graph>("tilelang_add_offline_graph");
	47	+ graph->SetInputs(inputs).SetOutputs(outputs);
	48	+ return graph;
	49	+}
	50	+} // namespace
	51	+
	52	+int main(int argc, char *argv[]) {
	53	+ (void)argc;
	54	+ (void)argv;
	55	+
	56	+ std::string output_air = "tilelang_add_offline.air";
	57	+ if (argc > 1) {
	58	+ output_air = argv[1];
	59	+ }
	60	+
	61	+ std::map<ge::AscendString, ge::AscendString> options = {
	62	+ {"ge.exec.deviceId", "0"},
	63	+ {"ge.graphRunMode", "1"},
	64	+ };
	65	+
	66	+ auto init_ret = ge::GEInitialize(options);
	67	+ if (init_ret != ge::SUCCESS) {
	68	+ std::cerr << "GEInitialize failed, ret: " << init_ret << std::endl;
	69	+ return 1;
	70	+ }
	71	+
	72	+ auto graph = BuildGraph();
	73	+
	74	+ std::cout << "Saving AIR file (for ATC offline compilation)..." << std::endl;
	75	+ auto ret = graph->SaveToFile(output_air);
	76	+ if (ret != ge::GRAPH_SUCCESS) {
	77	+ std::cerr << "SaveToFile failed, ret: " << ret << std::endl;
	78	+ (void)ge::GEFinalize();
	79	+ return 1;
	80	+ }
	81	+
	82	+ std::cout << "AIR file saved to: " << output_air << std::endl;
	83	+ std::cout << "Next step: run ATC to compile AIR to OM" << std::endl;
	84	+ std::cout << " atc --framework=1 --model=" << output_air << " --output=tilelang_add_offline"
	85	+ << " --soc_version=<your_soc>" << std::endl;
	86	+ std::cout << " (ATC will trigger Compile + Serialize)" << std::endl;
	87	+
	88	+ (void)ge::GEFinalize();
	89	+ return 0;
	90	+}

Aexamples/custom_op/tilelang_add_custom_offline/model_exec/main.cc+157-0

@@ -0,0 +1,157 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#include <cmath>
	12	+#include <cstring>
	13	+#include <iostream>
	14	+#include <random>
	15	+#include <string>
	16	+#include <vector>
	17	+
	18	+#include "acl/acl.h"
	19	+#include "acl/acl_rt.h"
	20	+
	21	+namespace {
	22	+constexpr int64_t kNumElements = 4096;
	23	+constexpr size_t kDataSizeBytes = static_cast<size_t>(kNumElements) * sizeof(float);
	24	+constexpr int kRandomSeed = 42;
	25	+constexpr float kErrorTolerance = 1e-5f;
	26	+constexpr int kMaxErrorDetails = 10;
	27	+
	28	+#define CHECK_ACL(ret, msg) \
	29	+ do { \
	30	+ if ((ret) != ACL_ERROR_NONE) { \
	31	+ std::cerr << (msg) << ", aclError: " << (ret) << std::endl; \
	32	+ return 1; \
	33	+ } \
	34	+ } while (0)
	35	+} // namespace
	36	+
	37	+int main(int argc, char *argv[]) {
	38	+ std::string om_path = "tilelang_add_offline.om";
	39	+ if (argc > 1) {
	40	+ om_path = argv[1];
	41	+ }
	42	+
	43	+ CHECK_ACL(aclInit(nullptr), "aclInit failed");
	44	+ CHECK_ACL(aclrtSetDevice(0), "aclrtSetDevice failed");
	45	+
	46	+ uint32_t model_id = 0;
	47	+ std::cout << "Loading OM model (triggers Deserialize): " << om_path << std::endl;
	48	+ auto ret = aclmdlLoadFromFile(om_path.c_str(), &model_id);
	49	+ if (ret != ACL_ERROR_NONE) {
	50	+ std::cerr << "aclmdlLoadFromFile failed, ret: " << ret << std::endl;
	51	+ CHECK_ACL(aclFinalize(), "aclFinalize failed");
	52	+ return 1;
	53	+ }
	54	+
	55	+ auto *model_desc = aclmdlCreateDesc();
	56	+ ret = aclmdlGetDesc(model_desc, model_id);
	57	+ if (ret != ACL_ERROR_NONE) {
	58	+ std::cerr << "aclmdlGetDesc failed, ret: " << ret << std::endl;
	59	+ (void)aclmdlUnload(model_id);
	60	+ CHECK_ACL(aclFinalize(), "aclFinalize failed");
	61	+ return 1;
	62	+ }
	63	+
	64	+ size_t input_size = aclmdlGetInputSizeByIndex(model_desc, 0);
	65	+ std::cout << "Input size: " << input_size << " bytes (expected " << kDataSizeBytes << ")" << std::endl;
	66	+ if (input_size != kDataSizeBytes) {
	67	+ std::cerr << "Input size mismatch!" << std::endl;
	68	+ (void)aclmdlUnload(model_id);
	69	+ (void)aclmdlDestroyDesc(model_desc);
	70	+ CHECK_ACL(aclFinalize(), "aclFinalize failed");
	71	+ return 1;
	72	+ }
	73	+
	74	+ size_t num_inputs = aclmdlGetNumInputs(model_desc);
	75	+ size_t num_outputs = aclmdlGetNumOutputs(model_desc);
	76	+ std::cout << "Model: " << num_inputs << " inputs, " << num_outputs << " outputs" << std::endl;
	77	+
	78	+ void *x_dev = nullptr;
	79	+ void *y_dev = nullptr;
	80	+ void *z_dev = nullptr;
	81	+ CHECK_ACL(aclrtMalloc(&x_dev, kDataSizeBytes, ACL_MEM_MALLOC_HUGE_FIRST), "aclrtMalloc x failed");
	82	+ CHECK_ACL(aclrtMalloc(&y_dev, kDataSizeBytes, ACL_MEM_MALLOC_HUGE_FIRST), "aclrtMalloc y failed");
	83	+ CHECK_ACL(aclrtMalloc(&z_dev, kDataSizeBytes, ACL_MEM_MALLOC_HUGE_FIRST), "aclrtMalloc z failed");
	84	+
	85	+ std::vector<float> host_x(kNumElements);
	86	+ std::vector<float> host_y(kNumElements);
	87	+ std::vector<float> host_z(kNumElements);
	88	+ std::mt19937 rng(kRandomSeed);
	89	+ std::uniform_real_distribution<float> dist(0.0f, 1.0f);
	90	+ for (int64_t i = 0; i < kNumElements; ++i) {
	91	+ host_x[i] = dist(rng);
	92	+ host_y[i] = dist(rng);
	93	+ }
	94	+ CHECK_ACL(aclrtMemcpy(x_dev, kDataSizeBytes, host_x.data(), kDataSizeBytes, ACL_MEMCPY_HOST_TO_DEVICE),
	95	+ "aclrtMemcpy x H2D failed");
	96	+ CHECK_ACL(aclrtMemcpy(y_dev, kDataSizeBytes, host_y.data(), kDataSizeBytes, ACL_MEMCPY_HOST_TO_DEVICE),
	97	+ "aclrtMemcpy y H2D failed");
	98	+
	99	+ aclmdlDataset *input_dataset = aclmdlCreateDataset();
	100	+ aclDataBuffer *x_buf = aclCreateDataBuffer(x_dev, kDataSizeBytes);
	101	+ aclDataBuffer *y_buf = aclCreateDataBuffer(y_dev, kDataSizeBytes);
	102	+ (void)aclmdlAddDatasetBuffer(input_dataset, x_buf);
	103	+ (void)aclmdlAddDatasetBuffer(input_dataset, y_buf);
	104	+
	105	+ aclmdlDataset *output_dataset = aclmdlCreateDataset();
	106	+ aclDataBuffer *z_buf = aclCreateDataBuffer(z_dev, kDataSizeBytes);
	107	+ (void)aclmdlAddDatasetBuffer(output_dataset, z_buf);
	108	+
	109	+ std::cout << "Executing model (triggers Execute)..." << std::endl;
	110	+ ret = aclmdlExecute(model_id, input_dataset, output_dataset);
	111	+ bool precision_ok = false;
	112	+ if (ret != ACL_ERROR_NONE) {
	113	+ std::cerr << "aclmdlExecute failed, ret: " << ret << std::endl;
	114	+ } else {
	115	+ CHECK_ACL(aclrtMemcpy(host_z.data(), kDataSizeBytes, z_dev, kDataSizeBytes, ACL_MEMCPY_DEVICE_TO_HOST),
	116	+ "aclrtMemcpy z D2H failed");
	117	+
	118	+ int error_count = 0;
	119	+ float max_error = 0.0f;
	120	+ for (int64_t i = 0; i < kNumElements; ++i) {
	121	+ float expected = host_x[i] + host_y[i];
	122	+ if (std::isnan(host_z[i]) \|\| std::isnan(expected)) {
	123	+ std::cerr << "NaN detected at [" << i << "]" << std::endl;
	124	+ error_count++;
	125	+ continue;
	126	+ }
	127	+ float error = std::abs(host_z[i] - expected);
	128	+ max_error = std::max(max_error, error);
	129	+ if (error > kErrorTolerance) {
	130	+ if (error_count < kMaxErrorDetails) {
	131	+ std::cerr << "Error at [" << i << "]: expected=" << expected << ", got=" << host_z[i] << std::endl;
	132	+ }
	133	+ error_count++;
	134	+ }
	135	+ }
	136	+ if (error_count > 0) {
	137	+ std::cerr << "Precision check failed: " << error_count << " errors, max_error=" << max_error << std::endl;
	138	+ } else {
	139	+ std::cout << "Precision check passed, max_error=" << max_error << std::endl;
	140	+ precision_ok = true;
	141	+ }
	142	+ }
	143	+
	144	+ (void)aclDestroyDataBuffer(x_buf);
	145	+ (void)aclDestroyDataBuffer(y_buf);
	146	+ (void)aclDestroyDataBuffer(z_buf);
	147	+ (void)aclmdlDestroyDataset(input_dataset);
	148	+ (void)aclmdlDestroyDataset(output_dataset);
	149	+ CHECK_ACL(aclrtFree(x_dev), "aclrtFree x failed");
	150	+ CHECK_ACL(aclrtFree(y_dev), "aclrtFree y failed");
	151	+ CHECK_ACL(aclrtFree(z_dev), "aclrtFree z failed");
	152	+
	153	+ (void)aclmdlUnload(model_id);
	154	+ (void)aclmdlDestroyDesc(model_desc);
	155	+ CHECK_ACL(aclFinalize(), "aclFinalize failed");
	156	+ return (ret == ACL_ERROR_NONE && precision_ok) ? 0 : 1;
	157	+}

Aexamples/custom_op/tilelang_add_custom_offline/run.sh+96-0

@@ -0,0 +1,96 @@
	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	+PROJECT_DIR="${SCRIPT_DIR}"
	16	+BUILD_DIR="${PROJECT_DIR}/build"
	17	+OUTPUT_DIR="${PROJECT_DIR}/output"
	18	+AIR_PATH="${BUILD_DIR}/tilelang_add_offline.air"
	19	+OM_PATH="${BUILD_DIR}/tilelang_add_offline"
	20	+
	21	+info() {
	22	+ echo "[INFO] $*"
	23	+}
	24	+
	25	+error() {
	26	+ echo "[ERROR] $*" >&2
	27	+}
	28	+
	29	+if [[ -z "${ASCEND_HOME_PATH:-}" ]]; then
	30	+ error "ASCEND_HOME_PATH is empty. Please source CANN set_env.sh first."
	31	+ exit 1
	32	+fi
	33	+
	34	+if [[ -n "${TILELANG_ASCEND_HOME:-}" ]]; then
	35	+ export PYTHONPATH="${TILELANG_ASCEND_HOME}:${PYTHONPATH:-}"
	36	+ export LD_LIBRARY_PATH="${TILELANG_ASCEND_HOME}/build:${LD_LIBRARY_PATH:-}"
	37	+ info "Using TILELANG_ASCEND_HOME=${TILELANG_ASCEND_HOME}"
	38	+else
	39	+ if ! python3 -c "import tilelang" 2>/dev/null; then
	40	+ error "tilelang is not importable. Set TILELANG_ASCEND_HOME or install tilelang-ascend."
	41	+ exit 1
	42	+ fi
	43	+fi
	44	+
	45	+SOC_VERSION="${SOC_VERSION:-Ascend910_9362}"
	46	+
	47	+mkdir -p "${BUILD_DIR}" "${OUTPUT_DIR}"
	48	+
	49	+info "Step 1/4: build custom op library, graph_build and model_exec"
	50	+cmake -S "${PROJECT_DIR}" -B "${BUILD_DIR}" -DCMAKE_BUILD_TYPE=Release
	51	+cmake --build "${BUILD_DIR}" -j"$(nproc 2>/dev/null \|\| echo 8)"
	52	+cmake --install "${BUILD_DIR}"
	53	+
	54	+export ASCEND_CUSTOM_OPP_PATH="${OUTPUT_DIR}:${ASCEND_CUSTOM_OPP_PATH:-}"
	55	+info "ASCEND_CUSTOM_OPP_PATH=${ASCEND_CUSTOM_OPP_PATH}"
	56	+
	57	+KERNEL_PY="${OUTPUT_DIR}/op_graph/lib/linux/$(uname -m \| tr '[:upper:]' '[:lower:]' \| sed 's/x86_64/x86_64/;s/aarch64/aarch64/')/add_custom_kernel.py"
	58	+if [[ ! -f "${KERNEL_PY}" ]]; then
	59	+ error "kernel source not found in OPP package: ${KERNEL_PY}"
	60	+ exit 1
	61	+fi
	62	+info "kernel source installed in OPP package."
	63	+
	64	+info "Step 2/4: generate AIR file (graph definition)"
	65	+"${BUILD_DIR}/tilelang_offline_graph_build" "${AIR_PATH}"
	66	+if [[ ! -f "${AIR_PATH}" ]]; then
	67	+ error "AIR file not generated: ${AIR_PATH}"
	68	+ exit 1
	69	+fi
	70	+info "AIR file generated: ${AIR_PATH}"
	71	+
	72	+info "Step 3/4: compile AIR to OM via ATC (triggers Compile + Serialize)"
	73	+if ! command -v atc &>/dev/null; then
	74	+ error "atc command not found. Please ensure CANN toolkit is properly installed."
	75	+ error "You can manually compile the AIR file:"
	76	+ error " atc --framework=1 --model=${AIR_PATH} --output=${OM_PATH} --soc_version=${SOC_VERSION}"
	77	+ exit 1
	78	+fi
	79	+
	80	+atc --framework=1 --model="${AIR_PATH}" --output="${OM_PATH}" --soc_version="${SOC_VERSION}" 2>&1 \|\| {
	81	+ error "ATC compilation failed. Check ATC logs for details."
	82	+ exit 1
	83	+}
	84	+
	85	+OM_FILE="${OM_PATH}.om"
	86	+if [[ ! -f "${OM_FILE}" ]]; then
	87	+ error "OM model not generated: ${OM_FILE}"
	88	+ exit 1
	89	+fi
	90	+info "OM model generated: $(ls -la "${OM_FILE}" \| awk '{print $5}') bytes"
	91	+
	92	+info "Step 4/4: execute OM model (triggers Deserialize + Execute)"
	93	+unset ASCEND_CUSTOM_OPP_PATH
	94	+"${BUILD_DIR}/tilelang_offline_model_exec" "${OM_FILE}"
	C C Chang-an-HW21 天前离线加载前应该先将ASCEND_CUSTOM_OPP_PATH环境变量清掉，加一行unset ASCEND_CUSTOM_OPP_PATH，跑一下看看有没有问题
	95	+
	96	+info "Sample pipeline finished."

Aexamples/custom_op/tilelang_add_custom_online/CMakeLists.txt+140-0

@@ -0,0 +1,140 @@
	1	+cmake_minimum_required(VERSION 3.14)
	2	+project(tilelang_add_custom_online LANGUAGES CXX)
	3	+
	4	+option(TILELANG_ONLINE_BUILD_CUSTOM_OP "Build libcust_opapi.so" ON)
	5	+option(TILELANG_ONLINE_BUILD_SESSION_RUN "Build session_run" ON)
	6	+
	7	+set(CMAKE_CXX_STANDARD 17)
	8	+set(CMAKE_CXX_STANDARD_REQUIRED ON)
	9	+set(CMAKE_CXX_EXTENSIONS OFF)
	10	+
	11	+if(NOT CMAKE_BUILD_TYPE)
	12	+ set(CMAKE_BUILD_TYPE Release CACHE STRING "Build type" FORCE)
	13	+endif()
	14	+
	15	+set(COMMON_COMPILE_OPTIONS
	16	+ -Wall
	17	+ -Wextra
	18	+ -Wno-unused-parameter
	19	+)
	20	+
	21	+set(PROJECT_OUTPUT_DIR "${CMAKE_SOURCE_DIR}/output")
	22	+file(MAKE_DIRECTORY "${PROJECT_OUTPUT_DIR}")
	23	+
	24	+if(CMAKE_SYSTEM_NAME STREQUAL "Windows")
	25	+ set(OPP_OS_TYPE "windows")
	26	+else()
	27	+ set(OPP_OS_TYPE "linux")
	28	+endif()
	29	+
	30	+string(TOLOWER "${CMAKE_SYSTEM_PROCESSOR}" CMAKE_SYSTEM_PROCESSOR_LOWER)
	31	+if(CMAKE_SYSTEM_PROCESSOR_LOWER MATCHES "^(aarch64\|arm64)$")
	32	+ set(OPP_CPU_TYPE "aarch64")
	33	+elseif(CMAKE_SYSTEM_PROCESSOR_LOWER MATCHES "^(x86_64\|amd64)$")
	34	+ set(OPP_CPU_TYPE "x86_64")
	35	+else()
	36	+ set(OPP_CPU_TYPE "${CMAKE_SYSTEM_PROCESSOR_LOWER}")
	37	+endif()
	38	+
	39	+set(CUSTOM_OP_OUTPUT_DIR "${PROJECT_OUTPUT_DIR}/op_graph/lib/${OPP_OS_TYPE}/${OPP_CPU_TYPE}")
	40	+file(MAKE_DIRECTORY "${CUSTOM_OP_OUTPUT_DIR}")
	41	+set(CUSTOM_OP_INCLUDE_DIR "${PROJECT_OUTPUT_DIR}/op_graph/include")
	42	+file(MAKE_DIRECTORY "${CUSTOM_OP_INCLUDE_DIR}")
	43	+
	44	+configure_file("${CMAKE_SOURCE_DIR}/ge/add_custom.h" "${CUSTOM_OP_INCLUDE_DIR}/add_custom.h" COPYONLY)
	45	+
	46	+set(KERNEL_SOURCE_FILE "${CMAKE_SOURCE_DIR}/add_custom_kernel/add_custom_kernel.py")
	47	+
	48	+set(ASCEND_HOME_PATH_OVERRIDE "" CACHE PATH "Optional ASCEND_HOME_PATH override")
	49	+if(ASCEND_HOME_PATH_OVERRIDE)
	50	+ set(ASCEND_HOME_PATH "${ASCEND_HOME_PATH_OVERRIDE}")
	51	+else()
	52	+ set(ASCEND_HOME_PATH "$ENV{ASCEND_HOME_PATH}")
	53	+endif()
	54	+
	55	+if(ASCEND_HOME_PATH)
	56	+ message(STATUS "ASCEND_HOME_PATH: ${ASCEND_HOME_PATH}")
	57	+else()
	58	+ message(WARNING "ASCEND_HOME_PATH is empty. Configure succeeds, but compilation requires a valid CANN toolkit path.")
	59	+endif()
	60	+
	61	+if(TILELANG_ONLINE_BUILD_CUSTOM_OP)
	62	+ configure_file("${KERNEL_SOURCE_FILE}" "${CUSTOM_OP_OUTPUT_DIR}/add_custom_kernel.py" COPYONLY)
	63	+
	64	+ add_library(cust_opapi SHARED
	65	+ ge/custom_op.cpp
	66	+ )
	67	+ target_compile_options(cust_opapi PRIVATE
	68	+ ${COMMON_COMPILE_OPTIONS}
	69	+ )
	70	+ target_compile_definitions(cust_opapi PRIVATE
	71	+ _GLIBCXX_USE_CXX11_ABI=0
	72	+ )
	73	+
	74	+ if(ASCEND_HOME_PATH)
	75	+ target_include_directories(cust_opapi PRIVATE
	76	+ "${CMAKE_SOURCE_DIR}/ge"
	77	+ "${ASCEND_HOME_PATH}/include"
	78	+ "${ASCEND_HOME_PATH}/include/graph"
	79	+ "${ASCEND_HOME_PATH}/include/register"
	80	+ "${ASCEND_HOME_PATH}/include/external"
	81	+ )
	82	+ target_link_directories(cust_opapi PRIVATE
	83	+ "${ASCEND_HOME_PATH}/lib64"
	84	+ )
	85	+ target_link_libraries(cust_opapi PRIVATE
	86	+ -Wl,--no-as-needed
	87	+ ascendcl
	88	+ lowering
	89	+ dl
	90	+ -Wl,--as-needed
	91	+ )
	92	+ endif()
	93	+
	94	+ set_target_properties(cust_opapi PROPERTIES
	95	+ OUTPUT_NAME "cust_opapi"
	96	+ LIBRARY_OUTPUT_DIRECTORY "${CUSTOM_OP_OUTPUT_DIR}"
	97	+ RUNTIME_OUTPUT_DIRECTORY "${CUSTOM_OP_OUTPUT_DIR}"
	98	+ ARCHIVE_OUTPUT_DIRECTORY "${CUSTOM_OP_OUTPUT_DIR}"
	99	+ )
	100	+
	101	+ install(FILES "${CUSTOM_OP_OUTPUT_DIR}/add_custom_kernel.py"
	102	+ DESTINATION "${CUSTOM_OP_OUTPUT_DIR}"
	103	+ )
	104	+ install(FILES "${CUSTOM_OP_INCLUDE_DIR}/add_custom.h"
	105	+ DESTINATION "${CUSTOM_OP_INCLUDE_DIR}"
	106	+ )
	107	+endif()
	108	+
	109	+if(TILELANG_ONLINE_BUILD_SESSION_RUN)
	110	+ add_executable(tilelang_online_session_run
	111	+ session_run/main.cc
	112	+ )
	113	+ target_compile_options(tilelang_online_session_run PRIVATE ${COMMON_COMPILE_OPTIONS})
	114	+ target_compile_definitions(tilelang_online_session_run PRIVATE
	115	+ _GLIBCXX_USE_CXX11_ABI=0
	116	+ )
	117	+
	118	+ if(ASCEND_HOME_PATH)
	119	+ target_include_directories(tilelang_online_session_run PRIVATE
	120	+ "${CUSTOM_OP_INCLUDE_DIR}"
	121	+ "${CMAKE_SOURCE_DIR}/ge"
	122	+ "${ASCEND_HOME_PATH}/include"
	123	+ "${ASCEND_HOME_PATH}/include/graph"
	124	+ "${ASCEND_HOME_PATH}/include/ge"
	125	+ "${ASCEND_HOME_PATH}/opp/built-in/op_proto/inc"
	126	+ )
	127	+ target_link_directories(tilelang_online_session_run PRIVATE "${ASCEND_HOME_PATH}/lib64")
	128	+ target_link_libraries(tilelang_online_session_run PRIVATE
	129	+ graph
	130	+ ge_runner
	131	+ ge_compiler
	132	+ ascendcl
	133	+ graph_base
	134	+ )
	135	+ endif()
	136	+
	137	+ set_target_properties(tilelang_online_session_run PROPERTIES
	138	+ RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}"
	139	+ )
	140	+endif()

Aexamples/custom_op/tilelang_add_custom_online/README.md+208-0

@@ -0,0 +1,208 @@
	1	+# TileLang Add Custom Operator 在线编译样例
	2	+
	3	+## 样例概述
	4	+
	5	+- 构图入口: GE 原生 (Session API)
	6	+- 算子编程语言: TileLang
	7	+- 编译方式: GE 编译阶段通过 `CompilableOp::Compile` 回调 subprocess 调用 TileLang Python 编译器，在线编译 kernel 源码为 `.so`
	8	+- 核心链路: `TileLang kernel 源码 → GE Compile 回调 → subprocess 编译 → dlopen 加载 → Execute 执行`
	9	+- 场景: 场景 B — 在线编译 + 在线执行（`CompilableOp` + `EagerExecuteOp` + `ShapeInferOp`）
	10	+
	11	+本样例以 element-wise Add 算子为例，展示如何通过 `CompilableOp` 接口在 GE 编译阶段（`CompileGraph`）在线编译 TileLang kernel 源码，而非预先编译好 `.so` 再加载。与 [tilelang_add_custom](../tilelang_add_custom/README.md)（eager 模式）样例形成对比。
	12	+
	13	+## 与 eager 样例的区别
	14	+
	15	+\| 维度 \| eager 样例 (`tilelang_add_custom`) \| 在线编译样例 (本样例) \|
	16	+\|------\|-----------------------------------\|---------------------\|
	17	+\| 接口组合 \| `EagerExecuteOp` + `ShapeInferOp` \| `CompilableOp` + `EagerExecuteOp` + `ShapeInferOp` \|
	18	+\| 编译时机 \| `run.sh` 预先编译 `.so` \| `CompileGraph` 阶段在线编译 \|
	19	+\| 加载时机 \| `Execute` 首次调用时 lazy `dlopen` \| `Compile` 回调中 `dlopen`，`Execute` 直接使用缓存 \|
	20	+\| 编译触发 \| 人工运行 `python3 add_custom_kernel.py` \| GE `CustomGraphOptimizer` 回调 `Compile` \|
	21	+\| shape 缓存 \| 无（固定 N=4096） \| 按元素数量构建 key 缓存，支持多元素数量 \|
	22	+\| 线程安全 \| `std::once_flag` \| `std::mutex`（`Compile` 可能被并行调用） \|
	23	+
	24	+## 目录结构
	25	+
	26	+```text
	27	+tilelang_add_custom_online/
	28	+├── README.md
	29	+├── README_en.md
	30	+├── CMakeLists.txt # 构建 libcust_opapi.so + session_run + 安装 .py 源码
	31	+├── run.sh # 一键编译运行（不预编译 kernel）
	32	+├── add_custom_kernel/
	33	+│ └── add_custom_kernel.py # TileLang kernel 源码（接受 N 和输出路径参数）
	34	+├── ge/
	35	+│ ├── add_custom.h # REG_OP 算子 proto 定义
	36	+│ └── custom_op.cpp # CompilableOp + EagerExecuteOp + ShapeInferOp 实现
	37	+└── session_run/
	38	+ └── main.cc # GE 原生构图 + CompileGraph(触发在线编译) + 执行 + 精度校验
	39	+```
	40	+
	41	+## 核心流程
	42	+
	43	+```text
	44	+GE 编译阶段 (CompileGraph):
	45	+ CustomGraphOptimizer 回调 Compile(ctx)
	46	+ ├─ 读取输入元素数量 → 构建 binary key
	47	+ ├─ 若 key 未缓存:
	48	+ │ ├─ 定位 add_custom_kernel.py（OPP 包中，与 libcust_opapi.so 同目录）
	49	+ │ ├─ exec("python3 add_custom_kernel.py <N> <output.so>")（同机有卡编译）
	50	+ │ ├─ TileLang 编译器编译 kernel 源码 → 产出 .so（host-wrapper）
	51	+ │ └─ dlopen .so + dlsym("call") → 缓存函数指针（临时文件读取后立即 unlink）
	52	+ └─ 返回 GRAPH_SUCCESS
	53	+
	54	+GE 执行阶段 (ExecuteGraphWithStreamAsync):
	55	+ 回调 Execute(ctx)
	56	+ ├─ 读取输入元素数量 → 构建 binary key
	57	+ ├─ 从缓存获取 call 函数指针
	58	+ ├─ 分配输出 Tensor
	59	+ └─ call(x_ptr, y_ptr, z_ptr, stream) → NPU 执行
	60	+```
	61	+
	62	+TileLang-Ascend 编译后的 `.so` 导出函数签名为：
	63	+
	64	+```c
	65	+extern "C" void call(uint8_t* A_handle, uint8_t* B_handle, uint8_t* C_handle, aclrtStream stream)
	66	+```
	67	+
	68	+该函数内部封装了 `main_kernel<<<>>>` 的 launch 逻辑（含硬件调度地址获取、tiling 等），GE 侧无需手动拼装 args。
	69	+
	70	+## 前置依赖
	71	+
	72	+### CANN
	73	+
	74	+- 已正确安装并配置 CANN 环境（`source ${ASCEND_HOME_PATH}/set_env.sh`）
	75	+- 当前环境具备 ACL、GE、Graph 相关头文件与库
	76	+
	77	+### TileLang-Ascend
	78	+
	79	+需安装 TileLang 主包和 TileLang-Ascend 后端：
	80	+
	81	+```bash
	82	+pip install tilelang
	83	+# TileLang-Ascend 后端：从 https://github.com/tile-ai/tilelang-ascend 安装
	84	+```
	85	+
	86	+若 TileLang-Ascend 以源码方式安装（未 `pip install`），需设置环境变量：
	87	+
	88	+```bash
	89	+export TILELANG_ASCEND_HOME=/path/to/tilelang-ascend
	90	+```
	91	+
	92	+### 环境变量
	93	+
	94	+\| 变量 \| 必需 \| 说明 \|
	95	+\|------\|------\|------\|
	96	+\| `ASCEND_HOME_PATH` \| 是 \| CANN toolkit 路径 \|
	97	+\| `TILELANG_ASCEND_HOME` \| 否 \| TileLang-Ascend 源码安装路径（pip 安装则无需设置） \|
	98	+\| `ASCEND_CUSTOM_OPP_PATH` \| 自动 \| 由 `run.sh` 自动设置 \|
	99	+
	100	+## 快速运行
	101	+
	102	+```bash
	103	+source ${ASCEND_HOME_PATH}/set_env.sh
	104	+bash run.sh
	105	+```
	106	+
	107	+`run.sh` 依次执行 3 个步骤：
	108	+
	109	+1. 构建 `libcust_opapi.so` 和 `tilelang_online_session_run`，将 `add_custom_kernel.py` 安装到 OPP 包
	110	+2. 确认 kernel 源码已在 OPP 包中
	111	+3. 运行测试程序（`CompileGraph` 触发 TileLang 在线编译，然后执行并校验精度）
	112	+
	113	+> 注意：与 eager 样例不同，本样例不在 `run.sh` 中预编译 TileLang kernel。kernel 编译发生在 `session_run` 调用 `CompileGraph` 时，由 GE 回调 `CompilableOp::Compile` 触发。
	114	+
	115	+成功时终端输出：
	116	+
	117	+```text
	118	+[INFO] Step 1/3: build custom op library and session_run
	119	+...
	120	+[INFO] Step 2/3: run session test (CompileGraph triggers TileLang online compilation)
	121	+CompileGraph (triggers TileLang online compilation)...
	122	+Compiling TileLang kernel: python3 ".../add_custom_kernel.py" 4096 ".../tilelang_add_custom_online_4096.so" 2>&1
	123	+Kernel .so saved to: ...
	124	+TileLang kernel compiled and loaded, key=4096, so=...
	125	+Precision check passed, max_error=0
	126	+[INFO] Step 3/3: sample pipeline finished.
	127	+```
	128	+
	129	+## 关键文件说明
	130	+
	131	+### `ge/custom_op.cpp`
	132	+
	133	+GE 交付件，实现 `CompilableOp` + `EagerExecuteOp` + `ShapeInferOp`：
	134	+
	135	+- Compile:
	136	+ 1. 从 `ctx->GetInputTensor(0)` 读取输入 shape size，构建 binary key
	137	+ 2. 加锁检查缓存，若 key 已存在则直接返回（支持多 shape）
	138	+ 3. 通过 `dladdr` 定位 `libcust_opapi.so` 所在目录，找到 `add_custom_kernel.py`
	139	+ 4. `popen` 调用 `python3 add_custom_kernel.py <N> <output.so>` 编译 TileLang kernel
	140	+ 5. `dlopen` 编译产出的 `.so`，`dlsym("call")` 获取函数指针，缓存到 `kernel_entries_`
	141	+- Execute:
	142	+ 1. 读取输入 shape size，构建 key
	143	+ 2. 从 `kernel_entries_` 获取 `Compile` 阶段缓存的函数指针
	144	+ 3. 分配输出 Tensor，调用 `call(x_ptr, y_ptr, z_ptr, stream)`
	145	+- InferShape / InferDataType: 输出 shape 和 dtype 与输入相同
	146	+- 使用 `std::mutex` 保证线程安全（`CustomGraphOptimizer` 可能并行调用 `Compile`）
	147	+
	148	+### `ge/add_custom.h`
	149	+
	150	+`REG_OP(AddCustomOnline)` 声明算子的输入输出规格，供 GE 原生构图创建节点。
	151	+
	152	+### `add_custom_kernel/add_custom_kernel.py`
	153	+
	154	+TileLang kernel 源码，接受命令行参数：
	155	+
	156	+- 第 1 个参数：`N`（元素总数，默认 4096，需为 BLOCK_SIZE=1024 的整数倍）
	157	+- 第 2 个参数：`output_path`（产出 `.so` 的路径）
	158	+
	159	+### `session_run/main.cc`
	160	+
	161	+GE 原生构图测试程序：
	162	+
	163	+1. `GEInitialize` + 创建 `Session`
	164	+2. 构建 `Data → AddCustomOnline` 计算图
	165	+3. `AddGraph` → `CompileGraph`（触发 `CompilableOp::Compile` → TileLang 在线编译）
	166	+4. `LoadGraph`
	167	+5. 分配 device 内存，H2D 拷贝输入数据
	168	+6. `ExecuteGraphWithStreamAsync` 执行
	169	+7. D2H 拷贝输出，逐元素精度校验（含 NaN 检查）
	170	+
	171	+## 算子规格
	172	+
	173	+\| 项目 \| 内容 \|
	174	+\|------\|------\|
	175	+\| 算子类型 \| `AddCustomOnline` \|
	176	+\| 输入 \| `x` (float32), `y` (float32) \|
	177	+\| 输出 \| `z` (float32) \|
	178	+\| 输入 shape \| `[4096]` (固定) \|
	179	+\| 输出 shape \| `[4096]` \|
	180	+\| 格式 \| ND \|
	181	+\| kernel 名称 \| `main_kernel`（由 `call` 封装） \|
	182	+\| BLOCK_SIZE \| 1024 \|
	183	+
	184	+## 分步运行
	185	+
	186	+```bash
	187	+# 1. 构建（含将 add_custom_kernel.py 安装到 OPP 包）
	188	+cmake -S . -B build -DCMAKE_BUILD_TYPE=Release
	189	+cmake --build build -j$(nproc)
	190	+cmake --install build
	191	+
	192	+# 2. 配置环境变量
	193	+export ASCEND_CUSTOM_OPP_PATH="$(pwd)/output:$ASCEND_CUSTOM_OPP_PATH"
	194	+
	195	+# 3. 运行（CompileGraph 触发在线编译）
	196	+./build/tilelang_online_session_run
	197	+```
	198	+
	199	+## 注意事项
	200	+
	201	+- 同机有卡编译限定：TileLang-Ascend 当前通过 `torch.npu.get_device_name()` 做运行时平台检测，不支持离线指定目标架构。本样例仅适用于"编译机与目标机为同一 NPU"的场景。
	202	+- kernel 源码 `.py` 安装在 OPP 包的 `op_graph/lib/<os>/<arch>/` 目录下，与 `libcust_opapi.so` 同目录，`Compile` 通过 `dladdr` 定位。
	203	+- 编译产出的 `.so` 使用 `mkstemps` 生成唯一临时文件，读取后立即 `unlink`，不会残留。
	204	+- `ge.graphRunMode=1` 确保走在线执行链路（PRIORITY_GRAPH 模式）。
	205	+- `CompileGraph` 必须在 `ExecuteGraphWithStreamAsync` 之前调用，否则 `Execute` 找不到已编译的 kernel。
	206	+- 当前样例仅支持 float32，如需支持更多数据类型需调整 `REG_OP` 的 `DATATYPE` 约束和 TileLang kernel 的 dtype 参数。
	207	+- TileLang-Ascend 的平台检测基于 `torch.npu.get_device_name()`，Ascend910 映射为 A2 平台。
	208	+- 在线编译需要运行环境具备 Python + TileLang，适合开发阶段；生产部署可改用 eager 样例的预编译方式。

Aexamples/custom_op/tilelang_add_custom_online/README_en.md+115-0

@@ -0,0 +1,115 @@
	1	+# TileLang Add Custom Operator Online Compilation Sample
	2	+
	3	+## Sample Overview
	4	+
	5	+- Graph construction entry: GE native (Session API)
	6	+- Operator programming language: TileLang
	7	+- Compilation method: GE compile phase invokes TileLang Python compiler via `CompilableOp::Compile` callback (subprocess), compiling kernel source to `.so` online
	8	+- Core pipeline: `TileLang kernel source → GE Compile callback → subprocess compilation → dlopen load → Execute`
	9	+- Scenario: Scenario B — online compilation + online execution (`CompilableOp` + `EagerExecuteOp` + `ShapeInferOp`)
	10	+
	11	+This sample demonstrates how to compile TileLang kernel source online during GE's compile phase (`CompileGraph`) via the `CompilableOp` interface, rather than pre-compiling the `.so`. Contrast with the [tilelang_add_custom](../tilelang_add_custom/README.md) (eager mode) sample.
	12	+
	13	+## Differences from Eager Sample
	14	+
	15	+\| Dimension \| Eager (`tilelang_add_custom`) \| Online Compilation (this sample) \|
	16	+\|-----------\|-------------------------------\|----------------------------------\|
	17	+\| Interface combo \| `EagerExecuteOp` + `ShapeInferOp` \| `CompilableOp` + `EagerExecuteOp` + `ShapeInferOp` \|
	18	+\| Compilation timing \| Pre-compiled by `run.sh` \| Online during `CompileGraph` \|
	19	+\| Load timing \| Lazy `dlopen` at first `Execute` \| `dlopen` in `Compile`, cached for `Execute` \|
	20	+\| Compilation trigger \| Manual `python3 add_custom_kernel.py` \| GE `CustomGraphOptimizer` calls `Compile` \|
	21	+\| Shape caching \| None (fixed N=4096) \| Keyed by element count, supports multiple element counts \|
	22	+\| Thread safety \| `std::once_flag` \| `std::mutex` (`Compile` may be called in parallel) \|
	23	+
	24	+## Directory Structure
	25	+
	26	+```text
	27	+tilelang_add_custom_online/
	28	+├── README.md
	29	+├── README_en.md
	30	+├── CMakeLists.txt
	31	+├── run.sh
	32	+├── add_custom_kernel/
	33	+│ └── add_custom_kernel.py
	34	+├── ge/
	35	+│ ├── add_custom.h
	36	+│ └── custom_op.cpp
	37	+└── session_run/
	38	+ └── main.cc
	39	+```
	40	+
	41	+## Core Pipeline
	42	+
	43	+```text
	44	+GE compile phase (CompileGraph):
	45	+ CustomGraphOptimizer calls Compile(ctx)
	46	+ ├─ Read input shape → build binary key
	47	+ ├─ If key not cached:
	48	+ │ ├─ Locate add_custom_kernel.py (in OPP package, same dir as libcust_opapi.so)
	49	+ │ ├─ popen("python3 add_custom_kernel.py <N> <output.so>")
	50	+ │ ├─ TileLang compiler compiles kernel source → .so (host-wrapper)
	51	+ │ └─ dlopen .so + dlsym("call") → cache function pointer
	52	+ └─ Return GRAPH_SUCCESS
	53	+
	54	+GE execution phase (ExecuteGraphWithStreamAsync):
	55	+ Execute(ctx) called
	56	+ ├─ Read input shape → build binary key
	57	+ ├─ Get cached call function pointer
	58	+ ├─ Allocate output Tensor
	59	+ └─ call(x_ptr, y_ptr, z_ptr, stream) → NPU execution
	60	+```
	61	+
	62	+## Prerequisites
	63	+
	64	+### CANN
	65	+
	66	+- CANN environment properly installed and configured (`source ${ASCEND_HOME_PATH}/set_env.sh`)
	67	+
	68	+### TileLang-Ascend
	69	+
	70	+```bash
	71	+pip install tilelang
	72	+# TileLang-Ascend backend: install from https://github.com/tile-ai/tilelang-ascend
	73	+```
	74	+
	75	+If installed from source, set:
	76	+
	77	+```bash
	78	+export TILELANG_ASCEND_HOME=/path/to/tilelang-ascend
	79	+```
	80	+
	81	+## Quick Start
	82	+
	83	+```bash
	84	+source ${ASCEND_HOME_PATH}/set_env.sh
	85	+bash run.sh
	86	+```
	87	+
	88	+`run.sh` executes 3 steps:
	89	+
	90	+1. Build `libcust_opapi.so` and `tilelang_online_session_run`, install `add_custom_kernel.py` to OPP package
	91	+2. Verify kernel source is in OPP package
	92	+3. Run test program (`CompileGraph` triggers TileLang online compilation, then executes and verifies)
	93	+
	94	+> Note: Unlike the eager sample, this sample does NOT pre-compile the TileLang kernel in `run.sh`. Compilation happens when `session_run` calls `CompileGraph`, triggered by GE's `CompilableOp::Compile` callback.
	95	+
	96	+## Operator Specification
	97	+
	98	+\| Item \| Value \|
	99	+\|------\|-------\|
	100	+\| Op type \| `AddCustomOnline` \|
	101	+\| Inputs \| `x` (float32), `y` (float32) \|
	102	+\| Output \| `z` (float32) \|
	103	+\| Input shape \| `[4096]` (fixed) \|
	104	+\| Format \| ND \|
	105	+\| Kernel name \| `main_kernel` (wrapped by `call`) \|
	106	+\| BLOCK_SIZE \| 1024 \|
	107	+
	108	+## Notes
	109	+
	110	+- Same-machine NPU compilation required: TileLang-Ascend uses `torch.npu.get_device_name()` for runtime platform detection and does not support specifying target architecture offline.
	111	+- Kernel source `.py` is installed in the OPP package at `op_graph/lib/<os>/<arch>/`, alongside `libcust_opapi.so`. `Compile` locates it via `dladdr`.
	112	+- Compiled `.so` uses `mkstemps` for unique temp file, unlinked immediately after reading.
	113	+- `ge.graphRunMode=1` ensures online execution (PRIORITY_GRAPH mode).
	114	+- `CompileGraph` must be called before `ExecuteGraphWithStreamAsync`, otherwise `Execute` cannot find the compiled kernel.
	115	+- Online compilation requires Python + TileLang in the runtime environment, suitable for development; for production deployment, consider the eager sample's pre-compilation approach.

Aexamples/custom_op/tilelang_add_custom_online/add_custom_kernel/add_custom_kernel.py+105-0

@@ -0,0 +1,105 @@
	1	+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
	2	+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	3	+# CANN Open Software License Agreement Version 2.0 (the "License").
	4	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	5	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	6	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	7	+# See LICENSE in the root of the software repository for the full text of the License.
	8	+
	9	+"""
	10	+TileLang Add Kernel + 在线编译产出 .so
	11	+
	12	+本文件用 TileLang 实现 element-wise Add kernel，并编译产出 Ascend .so 交付件。
	13	+供 GE 自定义算子 (CompilableOp::Compile) 在 GE 编译阶段通过 subprocess 调用本脚本完成在线编译。
	14	+
	15	+用法:
	16	+ python3 add_custom_kernel.py [N] [output_path]
	17	+
	18	+参数:
	19	+ N - 元素总数（默认 4096，需为 BLOCK_SIZE 的整数倍）
	20	+ output_path - 产出 .so 的路径（默认脚本目录下 add_kernel.so）
	21	+
	22	+TileLang-Ascend 编译后的 .so 导出函数签名为:
	23	+ extern "C" void call(uint8_t* A_handle, uint8_t* B_handle, uint8_t* C_handle, aclrtStream stream)
	24	+内部封装了 main_kernel<<<>>> 的 launch 逻辑。
	25	+"""
	26	+
	27	+import os
	28	+import shutil
	29	+import sys
	30	+
	31	+import tilelang
	32	+import tilelang.language as T
	33	+
	34	+BLOCK_SIZE = 1024
	35	+
	36	+
	37	+@tilelang.jit(out_idx=[-1])
	38	+def vec_add(n, block_size, dtype="float"):
	39	+ m_num = n // block_size
	40	+ vec_num = 2
	41	+
	42	+ @T.prim_func
	43	+ def main(
	44	+ a: T.Tensor((n,), dtype),
	45	+ b: T.Tensor((n,), dtype),
	46	+ c: T.Tensor((n,), dtype),
	47	+ ):
	48	+ with T.Kernel(m_num, is_npu=True) as (cid, vid):
	49	+ a_ub = T.alloc_ub((block_size // vec_num,), dtype)
	50	+ b_ub = T.alloc_ub((block_size // vec_num,), dtype)
	51	+ c_ub = T.alloc_ub((block_size // vec_num,), dtype)
	52	+ with T.Scope("V"):
	53	+ T.copy(a[cid * block_size + vid * block_size // vec_num], a_ub)
	54	+ T.copy(b[cid * block_size + vid * block_size // vec_num], b_ub)
	55	+
	56	+ T.barrier_all()
	57	+ T.tile.add(c_ub, a_ub, b_ub)
	58	+ T.barrier_all()
	59	+
	60	+ T.copy(c_ub, c[cid * block_size + vid * block_size // vec_num])
	61	+
	62	+ return main
	63	+
	64	+
	65	+def main():
	66	+ n = int(sys.argv[1]) if len(sys.argv) > 1 else 4096
	67	+ output_path = sys.argv[2] if len(sys.argv) > 2 else None
	68	+
	69	+ if n % BLOCK_SIZE != 0:
	70	+ print(
	71	+ f"Error: N={n} must be a multiple of BLOCK_SIZE={BLOCK_SIZE}",
	72	+ file=sys.stderr,
	73	+ )
	74	+ sys.exit(1)
	75	+
	76	+ func = vec_add(n, BLOCK_SIZE)
	77	+
	78	+ adapter = func.adapter
	79	+ so_path = getattr(getattr(adapter, "lib", None), "_name", None)
	80	+ if so_path is None:
	81	+ print(
	82	+ "Error: cannot locate compiled .so path from tilelang adapter",
	83	+ file=sys.stderr,
	84	+ )
	85	+ sys.exit(1)
	86	+
	87	+ if output_path is None:
	88	+ output_path = os.path.join(
	89	+ os.path.dirname(os.path.abspath(__file__)), "add_kernel.so"
	90	+ )
	91	+
	92	+ output_dir = os.path.dirname(os.path.abspath(output_path))
	93	+ os.makedirs(output_dir, exist_ok=True)
	94	+ shutil.copy2(so_path, output_path)
	95	+
	96	+ print(f"Kernel .so saved to: {output_path}")
	97	+ print(f"N={n}, BLOCK_SIZE={BLOCK_SIZE}")
	98	+ print(f"File size: {os.path.getsize(output_path)} bytes")
	99	+ print(
	100	+ "Export function: call(uint8_t* A, uint8_t* B, uint8_t* C, aclrtStream stream)"
	101	+ )
	102	+
	103	+
	104	+if __name__ == "__main__":
	105	+ main()

Aexamples/custom_op/tilelang_add_custom_online/ge/add_custom.h+25-0

@@ -0,0 +1,25 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+#ifndef EXAMPLES_CUSTOM_OP_TILELANG_ADD_CUSTOM_ONLINE_GE_ADD_CUSTOM_H_
	12	+#define EXAMPLES_CUSTOM_OP_TILELANG_ADD_CUSTOM_ONLINE_GE_ADD_CUSTOM_H_
	13	+
	14	+#include "graph/operator_reg.h"
	15	+
	16	+namespace ge {
	17	+REG_OP(AddCustomOnline)
	18	+ .INPUT(x, "T")
	19	+ .INPUT(y, "T")
	20	+ .OUTPUT(z, "T")
	21	+ .DATATYPE(T, TensorType({DT_FLOAT}))
	22	+ .OP_END_FACTORY_REG(AddCustomOnline);
	23	+} // namespace ge
	24	+
	25	+#endif // EXAMPLES_CUSTOM_OP_TILELANG_ADD_CUSTOM_ONLINE_GE_ADD_CUSTOM_H_

Aexamples/custom_op/tilelang_add_custom_online/ge/custom_op.cpp+236-0

Aexamples/custom_op/tilelang_add_custom_online/run.sh+63-0

Aexamples/custom_op/tilelang_add_custom_online/session_run/main.cc+205-0