docs(custom_op): 补齐 tilelang 样例英文 README 与中文版对齐-ge-AtomGit

已合并

docs(custom_op): 补齐 tilelang 样例英文 README 与中文版对齐 #4853

why you创建于 29 天前从Dazhou_Cao/ge： docs/tilelang-readme-en-sync 合入到 cann/ge： develop

docs(custom_op): 补齐 tilelang 样例英文 README 与中文版对齐 #4853

已合并

why you创建于 29 天前

从Dazhou_Cao/ge docs/tilelang-readme-en-sync合入到cann/ge develop

共 6 个文件变更+411-62

Mexamples/custom_op/tilelang_add_custom/README.md+6-6

	@@ -8,7 +8,7 @@
8	- 核心链路: `TileLang kernel → 预编译 .so → GE 交付件 → 进程内构图 → Session::ExecuteGraphWithStreamAsync 在线执行`	8	- 核心链路: `TileLang kernel → 预编译 .so → GE 交付件 → 进程内构图 → Session::ExecuteGraphWithStreamAsync 在线执行`
9	- 场景: 场景 A — 动态图在线执行（预编译 kernel + host 调度）	9	- 场景: 场景 A — 动态图在线执行（预编译 kernel + host 调度）
10		10
11	-本样例以 element-wise Add 算子为例，展示如何将 TileLang 编写的 kernel 通过 GE 语言无关自定义算子机制接入图编译和执行流程。	11	+本样例以 element-wise Add 算子为例，展示如何将 TileLang 编写的 kernel 通过 GE 的语言无关自定义算子机制接入图编译和执行流程。
12		12
13	## 目录结构	13	## 目录结构
14		14
	@@ -98,7 +98,7 @@ bash run.sh
98	Kernel .so saved to: add_kernel.so	98	Kernel .so saved to: add_kernel.so
99	[INFO] Step 2/4: build custom op library and session_run	99	[INFO] Step 2/4: build custom op library and session_run
100	...	100	...
101	-[INFO] Step 3/4: kernel .so installed in OPP package.	101	+[INFO] kernel .so installed in OPP package.
102	[INFO] Step 4/4: run session test	102	[INFO] Step 4/4: run session test
103	Precision check passed, max_error=0	103	Precision check passed, max_error=0
104	[INFO] Sample pipeline finished.	104	[INFO] Sample pipeline finished.
	@@ -111,12 +111,12 @@ Precision check passed, max_error=0
111	GE 交付件，实现 `EagerExecuteOp` + `ShapeInferOp`：	111	GE 交付件，实现 `EagerExecuteOp` + `ShapeInferOp`：
112		112
113	- Execute:	113	- Execute:
114	- 1. 首次调用时通过 `dlopen` 加载 `add_kernel.so`（路径从 `ASCEND_CUSTOM_OPP_PATH` 定位），`dlsym` 获取 `call` 函数指针	114	+ 1. 首次调用时通过 `dlopen` 加载 `add_kernel.so`（路径通过 `ASCEND_CUSTOM_OPP_PATH` 环境变量定位），`dlsym` 获取 `call` 函数指针
115	2. 校验两个输入的 shape size 均为 4096	115	2. 校验两个输入的 shape size 均为 4096
116	3. 分配输出 Tensor，调用 `call(x_ptr, y_ptr, z_ptr, stream)`	116	3. 分配输出 Tensor，调用 `call(x_ptr, y_ptr, z_ptr, stream)`
117	- InferShape / InferDataType: 输出 shape 和 dtype 与输入相同	117	- InferShape / InferDataType: 输出 shape 和 dtype 与输入相同
118	- 使用 `std::once_flag` 保证线程安全的延迟加载	118	- 使用 `std::once_flag` 保证线程安全的延迟加载
119	-- kernel `.so` 路径从 `ASCEND_CUSTOM_OPP_PATH` 环境变量定位，不依赖工作目录	119	+- kernel `.so` 路径通过 `ASCEND_CUSTOM_OPP_PATH` 环境变量定位，不依赖工作目录
120		120
121	### `ge/add_custom.h`	121	### `ge/add_custom.h`
122		122
	@@ -144,7 +144,7 @@ GE 原生构图测试程序：
144	\| 算子类型 \| `AddCustom` \|	144	\| 算子类型 \| `AddCustom` \|
145	\| 输入 \| `x` (float32), `y` (float32) \|	145	\| 输入 \| `x` (float32), `y` (float32) \|
146	\| 输出 \| `z` (float32) \|	146	\| 输出 \| `z` (float32) \|
147	-\| 输入 shape \| `[4096]` (固定) \|	147	+\| 输入 shape \| `[4096]`（固定） \|
148	\| 输出 shape \| `[4096]` \|	148	\| 输出 shape \| `[4096]` \|
149	\| 格式 \| ND \|	149	\| 格式 \| ND \|
150	\| kernel 名称 \| `main_kernel`（由 `call` 封装） \|	150	\| kernel 名称 \| `main_kernel`（由 `call` 封装） \|
	@@ -174,4 +174,4 @@ export ASCEND_CUSTOM_OPP_PATH="$(pwd)/output:$ASCEND_CUSTOM_OPP_PATH"
174	- `ge.graphRunMode=1` 确保走在线执行链路（PRIORITY_GRAPH 模式）。	174	- `ge.graphRunMode=1` 确保走在线执行链路（PRIORITY_GRAPH 模式）。
175	- 当前样例仅支持 float32，如需支持更多数据类型需调整 `REG_OP` 的 `DATATYPE` 约束和 TileLang kernel 的 dtype 参数。	175	- 当前样例仅支持 float32，如需支持更多数据类型需调整 `REG_OP` 的 `DATATYPE` 约束和 TileLang kernel 的 dtype 参数。
176	- TileLang-Ascend 的平台检测基于 `torch.npu.get_device_name()`，Ascend910 映射为 A2 平台。	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` 定位。	177	+- kernel `.so` 安装在 OPP 包的 `op_graph/lib/<os>/<arch>/` 目录下，与 `libcust_opapi.so` 同目录，路径通过 `ASCEND_CUSTOM_OPP_PATH` 环境变量定位。

Mexamples/custom_op/tilelang_add_custom/README_en.md+103-12

	@@ -4,11 +4,11 @@
4		4
5	- Graph construction entry: GE native (Session API)	5	- Graph construction entry: GE native (Session API)
6	- Operator programming language: TileLang	6	- Operator programming language: TileLang
7	-- Compilation method: TileLang pre-compiles to host-wrapper `.so`, loaded via `dlopen` at runtime	7	+- Compilation method: TileLang pre-compiles a host-wrapper `.so`, loaded via `dlopen` by GE at runtime
8	-- Core pipeline: `TileLang kernel → pre-compiled .so → GE deliverable → in-process graph → Session::ExecuteGraphWithStreamAsync online execution`	8	+- Core pipeline: `TileLang kernel → pre-compiled .so → GE deliverable → in-process graph construction → Session::ExecuteGraphWithStreamAsync online execution`
9	-- Scenario: Scenario A — online execution with pre-compiled kernel	9	+- Scenario: Scenario A — dynamic graph online execution (pre-compiled kernel + host scheduling)
10		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.	11	+This sample uses an element-wise Add operator to demonstrate how to integrate a TileLang-written kernel into GE's graph compilation and execution flow via the language-independent custom operator mechanism.
12		12
13	## Directory Structure	13	## Directory Structure
14		14
	@@ -35,29 +35,48 @@ TileLang kernel source (add_custom_kernel.py)
35	add_kernel.so (host-wrapper, exports call function)	35	add_kernel.so (host-wrapper, exports call function)
36	↓ dlopen + dlsym("call")	36	↓ dlopen + dlsym("call")
37	GE custom operator (AddCustom, EagerExecuteOp)	37	GE custom operator (AddCustom, EagerExecuteOp)
38	- ↓ call(x_ptr, y_ptr, z_ptr, stream) — wraps main_kernel<<<>>> launch	38	+ ↓ call(x_ptr, y_ptr, z_ptr, stream) — wraps main_kernel<<<>>> launch internally
39	NPU execution	39	NPU execution
40	```	40	```
41		41
		42	+The `.so` compiled by TileLang-Ascend exports a function with the signature:
		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	+This function internally wraps the launch logic of `main_kernel<<<>>>` (including hardware scheduling address acquisition, tiling, etc.), so GE does not need to assemble args manually.
		49	+
42	## Prerequisites	50	## Prerequisites
43		51
44	### CANN	52	### CANN
45		53
46	- CANN environment properly installed and configured (`source ${ASCEND_HOME_PATH}/set_env.sh`)	54	- CANN environment properly installed and configured (`source ${ASCEND_HOME_PATH}/set_env.sh`)
		55	+- The environment provides ACL, GE, and Graph related headers and libraries
47		56
48	### TileLang-Ascend	57	### TileLang-Ascend
49		58
		59	+Install the TileLang main package and the TileLang-Ascend backend:
		60	+
50	```bash	61	```bash
51	-pip install tilelang	62	+pip install tilelang # main package
52	# TileLang-Ascend backend: install from https://github.com/tile-ai/tilelang-ascend	63	# TileLang-Ascend backend: install from https://github.com/tile-ai/tilelang-ascend
53	```	64	```
54		65
55	-If TileLang-Ascend is installed from source (not via pip), set:	66	+If TileLang-Ascend is installed from source (not via `pip install`), set the environment variable:
56		67
57	```bash	68	```bash
58	export TILELANG_ASCEND_HOME=/path/to/tilelang-ascend	69	export TILELANG_ASCEND_HOME=/path/to/tilelang-ascend
59	```	70	```
60		71
		72	+### Environment Variables
		73	+
		74	+\| Variable \| Required \| Description \|
		75	+\|----------\|----------\|-------------\|
		76	+\| `ASCEND_HOME_PATH` \| Yes \| CANN toolkit path \|
		77	+\| `TILELANG_ASCEND_HOME` \| No \| TileLang-Ascend source installation path (not needed if installed via pip) \|
		78	+\| `ASCEND_CUSTOM_OPP_PATH` \| Auto \| Set automatically by `run.sh` \|
		79	+
61	## Quick Start	80	## Quick Start
62		81
63	```bash	82	```bash
	@@ -65,6 +84,59 @@ source ${ASCEND_HOME_PATH}/set_env.sh
65	bash run.sh	84	bash run.sh
66	```	85	```
67		86
		87	+`run.sh` executes 4 steps in sequence:
		88	+
		89	+1. Compile the TileLang kernel, producing `add_kernel.so`
		90	+2. Build `libcust_opapi.so` and `tilelang_session_run`, install `add_kernel.so` into the OPP package
		91	+3. Verify the kernel `.so` is in the OPP package
		92	+4. Run the test program
		93	+
		94	+Expected terminal output on success:
		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] 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	+## Key Files
		108	+
		109	+### `ge/custom_op.cpp`
		110	+
		111	+GE deliverable implementing `EagerExecuteOp` + `ShapeInferOp`:
		112	+
		113	+- Execute:
		114	+ 1. On the first call, loads `add_kernel.so` via `dlopen` (path located via the `ASCEND_CUSTOM_OPP_PATH` environment variable) and obtains the `call` function pointer via `dlsym`
		115	+ 2. Validates that both inputs have a shape size of 4096
		116	+ 3. Allocates the output Tensor and calls `call(x_ptr, y_ptr, z_ptr, stream)`
		117	+- InferShape / InferDataType: output shape and dtype are the same as the input
		118	+- Uses `std::once_flag` for thread-safe lazy loading
		119	+- The kernel `.so` path is located via the `ASCEND_CUSTOM_OPP_PATH` environment variable, independent of the working directory
		120	+
		121	+### `ge/add_custom.h`
		122	+
		123	+`REG_OP(AddCustom)` declares the operator's input/output specification, used by GE native graph construction to create nodes.
		124	+
		125	+### `add_custom_kernel/add_custom_kernel.py`
		126	+
		127	+TileLang kernel source that defines the element-wise Add and compiles it into `add_kernel.so`.
		128	+
		129	+### `session_run/main.cc`
		130	+
		131	+GE native graph construction test program:
		132	+
		133	+1. `GEInitialize` + create a `Session`
		134	+2. Build the `Data → AddCustom` computation graph
		135	+3. `AddGraph` → `CompileGraph` → `LoadGraph`
		136	+4. Allocate device memory, H2D copy of input data
		137	+5. Execute with `ExecuteGraphWithStreamAsync`
		138	+6. D2H copy of the output, element-wise precision check (including NaN check)
		139	+
68	## Operator Specification	140	## Operator Specification
69		141
70	\| Item \| Value \|	142	\| Item \| Value \|
	@@ -73,14 +145,33 @@ bash run.sh
73	\| Inputs \| `x` (float32), `y` (float32) \|	145	\| Inputs \| `x` (float32), `y` (float32) \|
74	\| Output \| `z` (float32) \|	146	\| Output \| `z` (float32) \|
75	\| Input shape \| `[4096]` (fixed) \|	147	\| Input shape \| `[4096]` (fixed) \|
		148	+\| Output shape \| `[4096]` \|
76	\| Format \| ND \|	149	\| Format \| ND \|
77	\| Kernel name \| `main_kernel` (wrapped by `call`) \|	150	\| Kernel name \| `main_kernel` (wrapped by `call`) \|
78	\| BLOCK_SIZE \| 1024 \|	151	\| BLOCK_SIZE \| 1024 \|
79		152
		153	+## Step-by-Step Run
		154	+
		155	+```bash
		156	+# 1. Compile the TileLang kernel
		157	+cd add_custom_kernel && python3 add_custom_kernel.py && cd ..
		158	+
		159	+# 2. Build (including installing add_kernel.so into the OPP package)
		160	+cmake -S . -B build -DCMAKE_BUILD_TYPE=Release
		161	+cmake --build build -j$(nproc)
		162	+cmake --install build
		163	+
		164	+# 3. Set environment variables
		165	+export ASCEND_CUSTOM_OPP_PATH="$(pwd)/output:$ASCEND_CUSTOM_OPP_PATH"
		166	+
		167	+# 4. Run
		168	+./build/tilelang_session_run
		169	+```
		170	+
80	## Notes	171	## Notes
81		172
82	-- The kernel is compiled with fixed N=4096. Execute validates input shape size and returns failure on mismatch.	173	+- The kernel is compiled with fixed N=4096. Execute validates the input shape size and returns failure on mismatch.
83	-- `ge.graphRunMode=1` ensures online execution (PRIORITY_GRAPH mode).	174	+- `ge.graphRunMode=1` ensures the online execution path (PRIORITY_GRAPH mode).
84	-- Only float32 is supported. To support more data types, adjust `REG_OP` DATATYPE and TileLang kernel dtype parameter.	175	+- Only float32 is supported. To support more data types, adjust the `REG_OP` `DATATYPE` constraint and the TileLang kernel dtype parameter.
85	-- TileLang-Ascend platform detection is based on `torch.npu.get_device_name()`. Ascend910 maps to A2 platform.	176	+- TileLang-Ascend platform detection is based on `torch.npu.get_device_name()`. Ascend910 maps to the A2 platform.
86	-- The kernel `.so` is installed in the OPP package at `op_graph/lib/<os>/<arch>/`, alongside `libcust_opapi.so`.	177	+- The kernel `.so` is installed in the OPP package at `op_graph/lib/<os>/<arch>/`, alongside `libcust_opapi.so`. The path is located via the `ASCEND_CUSTOM_OPP_PATH` environment variable.

Mexamples/custom_op/tilelang_add_custom_offline/README.md+4-4

	@@ -48,7 +48,7 @@ graph_build 生成 AIR 文件 → ATC 加载 AIR 编译 OM
48		48
49	GE 回调 Compile(ctx)	49	GE 回调 Compile(ctx)
50	├─ 读取输入元素数量 → 构建 binary key	50	├─ 读取输入元素数量 → 构建 binary key
51	- ├─ exec("python3 add_custom_kernel.py <N> <output.so>")（同机有卡编译）	51	+ ├─ popen("python3 add_custom_kernel.py <N> <output.so>")（同机有卡编译）
52	├─ TileLang 编译器编译 kernel 源码 → 产出 .so（host-wrapper）	52	├─ TileLang 编译器编译 kernel 源码 → 产出 .so（host-wrapper）
53	├─ 读取 .so 文件字节 → so_data	53	├─ 读取 .so 文件字节 → so_data
54	├─ mkstemps 临时文件读取后立即 unlink	54	├─ mkstemps 临时文件读取后立即 unlink
	@@ -144,7 +144,7 @@ Precision check passed, max_error=0
144	\| 算子类型 \| `AddCustomOffline` \|	144	\| 算子类型 \| `AddCustomOffline` \|
145	\| 输入 \| `x` (float32), `y` (float32) \|	145	\| 输入 \| `x` (float32), `y` (float32) \|
146	\| 输出 \| `z` (float32) \|	146	\| 输出 \| `z` (float32) \|
147	-\| 输入 shape \| `[4096]` (固定) \|	147	+\| 输入 shape \| `[4096]`（固定） \|
148	\| 输出 shape \| `[4096]` \|	148	\| 输出 shape \| `[4096]` \|
149	\| 格式 \| ND \|	149	\| 格式 \| ND \|
150	\| kernel 名称 \| `main_kernel`（由 `call` 封装） \|	150	\| kernel 名称 \| `main_kernel`（由 `call` 封装） \|
	@@ -176,7 +176,7 @@ GE 交付件，实现 `CompilableOp` + `PortableOp` + `EagerExecuteOp` + `ShapeI
176		176
177	- Compile: subprocess 调用 Python 编译 TileLang → 读取 `.so` 字节 → `dlopen` → 缓存	177	- Compile: subprocess 调用 Python 编译 TileLang → 读取 `.so` 字节 → `dlopen` → 缓存
178	- Serialize: 将 `kernel_entries_` 中的 `.so` 字节序列化为二进制 buffer	178	- Serialize: 将 `kernel_entries_` 中的 `.so` 字节序列化为二进制 buffer
179	-- Deserialize: 从 buffer 恢复 `.so` 字节 → 写临时文件 → `dlopen` → 缓存	179	+- Deserialize: 从 buffer 恢复 `.so` 字节 → `memfd_create` 从内存加载（不落盘）→ `dlopen` → 缓存
180	- Execute: 使用缓存的函数指针调用 `call(x, y, z, stream)`	180	- Execute: 使用缓存的函数指针调用 `call(x, y, z, stream)`
181	- 使用 `std::mutex` 保证线程安全	181	- 使用 `std::mutex` 保证线程安全
182		182
	@@ -202,7 +202,7 @@ ATC 编译 AIR → OM 时会自动触发 `Compile` + `Serialize`。
202		202
203	## 注意事项	203	## 注意事项
204		204
205	-- 同机有卡编译限定：TileLang-Ascend 当前通过 `torch.npu.get_device_name()` 做运行时平台检测，不支持离线指定目标架构。因此 `Compile` 回调中调用 Python 编译器时未传递 `--soc_version`，编译产物绑定编译机 NPU。本样例仅适用于"编译机与目标机为同一 NPU"的场景，不能用于跨平台 ATC 离线编译。如需跨平台编译，需等待 TileLang 支持离线目标指定后更新。	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` 编译产物）。	206	- `graph_build` 阶段需要运行环境具备 Python + TileLang-Ascend；`model_exec` 阶段不需要（OM 自包含 `.so` 编译产物）。
207	- `ge.graphRunMode=1` 确保走在线执行链路（PRIORITY_GRAPH 模式）。	207	- `ge.graphRunMode=1` 确保走在线执行链路（PRIORITY_GRAPH 模式）。
208	- 序列化格式为自定义格式，GE 只透传不解析，格式完全由算子控制。所有 `uint32_t` 字段使用小端格式。	208	- 序列化格式为自定义格式，GE 只透传不解析，格式完全由算子控制。所有 `uint32_t` 字段使用小端格式。

Mexamples/custom_op/tilelang_add_custom_offline/README_en.md+173-8

	@@ -4,22 +4,103 @@
4		4
5	- Graph construction entry: GE native (`Graph::SaveToFile` generates AIR, then ATC compiles OM)	5	- Graph construction entry: GE native (`Graph::SaveToFile` generates AIR, then ATC compiles OM)
6	- Operator programming language: TileLang	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	7	+- Compilation method: ATC compile phase invokes the TileLang Python compiler via `CompilableOp::Compile` callback (subprocess), compiling kernel source to `.so` online, then `PortableOp::Serialize` serializes the `.so` bytes into the OM model
8	- Core pipeline: `Graph → AIR → ATC compile (Compile + Serialize) → OM → ACL load (Deserialize + Execute)`	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`)	9	- Scenario: Scenario C — offline OM model sinking (`CompilableOp` + `PortableOp` + `EagerExecuteOp` + `ShapeInferOp`)
10		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).	11	+This sample uses an element-wise Add operator to demonstrate how to serialize TileLang compilation products into an OM model file via the `PortableOp` interface, enabling an offline deployment path. Contrast with [tilelang_add_custom_online](../tilelang_add_custom_online/README_en.md) (online compilation, Scenario B).
12		12
13	## Differences from Online Compilation Sample	13	## Differences from Online Compilation Sample
14		14
15	-\| Dimension \| Online (`tilelang_add_custom_online`) \| Offline OM (this sample) \|	15	+\| Dimension \| Online (`tilelang_add_custom_online`) \| Offline OM sinking (this sample) \|
16	-\|-----------\|---------------------------------------\|--------------------------\|	16	+\|-----------\|---------------------------------------\|----------------------------------\|
17	\| Interface combo \| `CompilableOp` + `EagerExecuteOp` + `ShapeInferOp` \| + `PortableOp` \|	17	\| Interface combo \| `CompilableOp` + `EagerExecuteOp` + `ShapeInferOp` \| + `PortableOp` \|
18	\| Model format \| No OM, direct `Session::ExecuteGraphWithStreamAsync` \| OM model file \|	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 \|	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` \|	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 \|	21	\| Deployment \| Requires Python + TileLang in runtime \| OM file is self-contained, no Python + TileLang needed at deployment \|
22		22
		23	+## Directory Structure
		24	+
		25	+```text
		26	+tilelang_add_custom_offline/
		27	+├── README.md
		28	+├── README_en.md
		29	+├── CMakeLists.txt # Build libcust_opapi.so + graph_build + model_exec
		30	+├── run.sh # One-click build and run
		31	+├── add_custom_kernel/
		32	+│ └── add_custom_kernel.py # TileLang kernel source (accepts N and output path arguments)
		33	+├── ge/
		34	+│ ├── add_custom.h # REG_OP proto definition
		35	+│ └── custom_op.cpp # CompilableOp + PortableOp + EagerExecuteOp + ShapeInferOp implementation
		36	+├── graph_build/
		37	+│ └── main.cc # Graph construction + Graph::SaveToFile generates AIR (for ATC compilation)
		38	+└── model_exec/
		39	+ └── main.cc # ACL loads OM + execution + precision check (triggers Deserialize + Execute)
		40	+```
		41	+
		42	+## Core Pipeline
		43	+
		44	+```text
		45	+=== graph_build phase (Graph::SaveToFile → ATC) ===
		46	+
		47	+graph_build generates the AIR file → ATC loads AIR and compiles OM
		48	+
		49	+GE calls back Compile(ctx)
		50	+ ├─ Read input element count → build binary key
		51	+ ├─ popen("python3 add_custom_kernel.py <N> <output.so>") (same-machine NPU compilation)
		52	+ ├─ TileLang compiler compiles kernel source → produces .so (host-wrapper)
		53	+ ├─ Read .so file bytes → so_data
		54	+ ├─ mkstemps temp file unlinked immediately after reading
		55	+ └─ dlopen .so + dlsym("call") → cache function pointer
		56	+
		57	+GE calls back Serialize(buffer)
		58	+ ├─ Little-endian format: [magic][version][count]
		59	+ │ [key_len][key][so_size][so_data] ...
		60	+ └─ Write all .so bytes in kernel_entries_ into the buffer → embedded into OM
		61	+
		62	+aclgrphSaveModel → save the OM file
		63	+
		64	+=== model_exec phase (aclmdlLoadFromFile) ===
		65	+
		66	+ACL loads OM → GE calls back Deserialize(buffer)
		67	+ ├─ Validate magic/version/count, check boundaries and duplicate keys
		68	+ ├─ Restore kernel entries one by one, load .so from memory via memfd_create (no disk file)
		69	+ ├─ dlopen memfd + dlsym("call") → cache function pointer
		70	+ ├─ Check there is no trailing dirty data
		71	+ └─ Atomically replace kernel_entries_ after all entries succeed (transactional)
		72	+
		73	+aclmdlExecute → GE calls back Execute(ctx)
		74	+ ├─ Get the call function pointer from kernel_entries_
		75	+ ├─ Allocate output Tensor
		76	+ └─ call(x_ptr, y_ptr, z_ptr, stream) → NPU execution
		77	+```
		78	+
		79	+## Prerequisites
		80	+
		81	+### CANN
		82	+
		83	+- CANN environment properly installed and configured (`source ${ASCEND_HOME_PATH}/set_env.sh`)
		84	+
		85	+### TileLang-Ascend
		86	+
		87	+Install the TileLang main package and the TileLang-Ascend backend:
		88	+
		89	+```bash
		90	+pip install tilelang
		91	+# TileLang-Ascend backend: install from https://github.com/tile-ai/tilelang-ascend
		92	+```
		93	+
		94	+> Note: TileLang-Ascend is only needed in the graph_build phase (compiling OM); the model_exec phase (loading and executing OM) does not need it.
		95	+
		96	+### Environment Variables
		97	+
		98	+\| Variable \| Required \| Description \|
		99	+\|----------\|----------\|-------------\|
		100	+\| `ASCEND_HOME_PATH` \| Yes \| CANN toolkit path \|
		101	+\| `TILELANG_ASCEND_HOME` \| No \| TileLang-Ascend source installation path (not needed if installed via pip) \|
		102	+\| `ASCEND_CUSTOM_OPP_PATH` \| Auto \| Set automatically by `run.sh` \|
		103	+
23	## Quick Start	104	## Quick Start
24		105
25	```bash	106	```bash
	@@ -27,6 +108,35 @@ source ${ASCEND_HOME_PATH}/set_env.sh
27	bash run.sh	108	bash run.sh
28	```	109	```
29		110
		111	+`run.sh` executes 4 steps in sequence:
		112	+
		113	+1. Build `libcust_opapi.so`, `graph_build`, and `model_exec`, install the `.py` source into the OPP package
		114	+2. Run `graph_build` (`Graph::SaveToFile` generates the AIR file)
		115	+3. Run `atc` (compile AIR → OM, triggers `Compile` + `Serialize`)
		116	+4. Run `model_exec` (`aclmdlLoadFromFile` triggers `Deserialize`, `aclmdlExecute` triggers `Execute`)
		117	+
		118	+Expected terminal output on success:
		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	+
30	## Operator Specification	140	## Operator Specification
31		141
32	\| Item \| Value \|	142	\| Item \| Value \|
	@@ -35,13 +145,68 @@ bash run.sh
35	\| Inputs \| `x` (float32), `y` (float32) \|	145	\| Inputs \| `x` (float32), `y` (float32) \|
36	\| Output \| `z` (float32) \|	146	\| Output \| `z` (float32) \|
37	\| Input shape \| `[4096]` (fixed) \|	147	\| Input shape \| `[4096]` (fixed) \|
		148	+\| Output shape \| `[4096]` \|
38	\| Format \| ND \|	149	\| Format \| ND \|
39	\| Kernel name \| `main_kernel` (wrapped by `call`) \|	150	\| Kernel name \| `main_kernel` (wrapped by `call`) \|
40	\| BLOCK_SIZE \| 1024 \|	151	\| BLOCK_SIZE \| 1024 \|
41		152
		153	+## Serialization Format
		154	+
		155	+`PortableOp::Serialize` uses a custom binary format to embed the TileLang `.so` compilation product into OM:
		156	+
		157	+```text
		158	+Offset Length Field Description
		159	+0 4 magic Fixed 0x4F504B4E (custom format identifier, little-endian)
		160	+4 4 version Fixed 1 (little-endian)
		161	+8 4 count Number of kernel entries (little-endian)
		162	+12 --- entries Repeated count times:
		163	+ 4 key_len Key byte length (little-endian)
		164	+ N key Element count string (e.g. "4096")
		165	+ 4 so_size .so file byte length (little-endian)
		166	+ M so_data Full .so binary content
		167	+```
		168	+
		169	+`Deserialize` reads this format, loads each `.so` from memory via `memfd_create` (no disk file), and checks integrity constraints such as duplicate keys and trailing dirty data.
		170	+
		171	+## Key Files
		172	+
		173	+### `ge/custom_op.cpp`
		174	+
		175	+GE deliverable implementing `CompilableOp` + `PortableOp` + `EagerExecuteOp` + `ShapeInferOp`:
		176	+
		177	+- Compile: invokes Python via subprocess to compile TileLang → reads `.so` bytes → `dlopen` → cache
		178	+- Serialize: serializes the `.so` bytes in `kernel_entries_` into a binary buffer
		179	+- Deserialize: restores `.so` bytes from the buffer → loads from memory via `memfd_create` (no disk file) → `dlopen` → cache
		180	+- Execute: calls `call(x, y, z, stream)` with the cached function pointer
		181	+- Uses `std::mutex` for thread safety
		182	+
		183	+### `graph_build/main.cc`
		184	+
		185	+Generates the AIR file with `Graph::SaveToFile` for offline ATC compilation:
		186	+
		187	+1. `GEInitialize` + build the computation graph
		188	+2. `graph->SaveToFile(air_path)` — generate the AIR file
		189	+
		190	+When ATC compiles AIR → OM, `Compile` + `Serialize` are triggered automatically.
		191	+
		192	+### `model_exec/main.cc`
		193	+
		194	+Loads and executes the OM model with ACL APIs:
		195	+
		196	+1. `aclInit` + `aclrtSetDevice`
		197	+2. `aclmdlLoadFromFile(om_path)` — triggers Deserialize
		198	+3. `aclmdlGetDesc` to get the model description
		199	+4. Allocate device memory, H2D copy of inputs
		200	+5. `aclmdlExecute` — triggers Execute
		201	+6. D2H copy of the output, precision check
		202	+
42	## Notes	203	## Notes
43		204
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.	205	+- Same-machine NPU compilation required: TileLang-Ascend currently uses `torch.npu.get_device_name()` for runtime platform detection and does not support specifying the target architecture offline. Therefore the Python compiler invoked in the `Compile` callback is not passed `--soc_version`, and the compilation product is bound to the compilation machine's NPU. This sample only applies to scenarios where the compilation machine and the target machine have the same NPU; it cannot be used for cross-platform ATC offline compilation. For cross-platform compilation, wait for TileLang to support offline target specification, then update this sample.
45	-- `graph_build` phase requires Python + TileLang-Ascend; `model_exec` phase does not (OM is self-contained).	206	+- The `graph_build` phase requires Python + TileLang-Ascend in the runtime environment; the `model_exec` phase does not (OM is self-contained with the compiled `.so`).
46	-- Serialization format is custom (little-endian); GE only transparently passes the buffer.	207	+- `ge.graphRunMode=1` ensures the online execution path (PRIORITY_GRAPH mode).
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.	208	+- The serialization format is custom; GE only passes the buffer through without parsing, and the format is fully controlled by the operator. All `uint32_t` fields use little-endian format.
		209	+- Only float32 is supported. To support more data types, adjust the `REG_OP` `DATATYPE` constraint and the TileLang kernel dtype parameter.
		210	+- OM compilation depends on the ATC tool: `graph_build` generates the AIR file, and ATC compiles AIR → OM (triggering `Compile` + `Serialize`). The `SOC_VERSION` environment variable can override the default `Ascend910_9362`.
		211	+- The compiled `.so` uses `mkstemps` for a unique temp file during `Compile` and calls `unlink` immediately after reading; `Deserialize` loads it from memory via `memfd_create`, without touching disk.
		212	+- If ATC is unavailable in the current environment (version mismatch, etc.), refer to the `compilable_add_custom` sample to compile the OM in an environment where ATC is available.

Mexamples/custom_op/tilelang_add_custom_online/README.md+4-4

	@@ -18,7 +18,7 @@
18	\| 编译时机 \| `run.sh` 预先编译 `.so` \| `CompileGraph` 阶段在线编译 \|	18	\| 编译时机 \| `run.sh` 预先编译 `.so` \| `CompileGraph` 阶段在线编译 \|
19	\| 加载时机 \| `Execute` 首次调用时 lazy `dlopen` \| `Compile` 回调中 `dlopen`，`Execute` 直接使用缓存 \|	19	\| 加载时机 \| `Execute` 首次调用时 lazy `dlopen` \| `Compile` 回调中 `dlopen`，`Execute` 直接使用缓存 \|
20	\| 编译触发 \| 人工运行 `python3 add_custom_kernel.py` \| GE `CustomGraphOptimizer` 回调 `Compile` \|	20	\| 编译触发 \| 人工运行 `python3 add_custom_kernel.py` \| GE `CustomGraphOptimizer` 回调 `Compile` \|
21	-\| shape 缓存 \| 无（固定 N=4096） \| 按元素数量构建 key 缓存，支持多元素数量 \|	21	+\| shape 缓存 \| 无（固定 N=4096） \| 按元素数量构建 key 缓存，支持多种输入规模 \|
22	\| 线程安全 \| `std::once_flag` \| `std::mutex`（`Compile` 可能被并行调用） \|	22	\| 线程安全 \| `std::once_flag` \| `std::mutex`（`Compile` 可能被并行调用） \|
23		23
24	## 目录结构	24	## 目录结构
	@@ -46,7 +46,7 @@ GE 编译阶段 (CompileGraph):
46	├─ 读取输入元素数量 → 构建 binary key	46	├─ 读取输入元素数量 → 构建 binary key
47	├─ 若 key 未缓存:	47	├─ 若 key 未缓存:
48	│ ├─ 定位 add_custom_kernel.py（OPP 包中，与 libcust_opapi.so 同目录）	48	│ ├─ 定位 add_custom_kernel.py（OPP 包中，与 libcust_opapi.so 同目录）
49	- │ ├─ exec("python3 add_custom_kernel.py <N> <output.so>")（同机有卡编译）	49	+ │ ├─ popen("python3 add_custom_kernel.py <N> <output.so>")（同机有卡编译）
50	│ ├─ TileLang 编译器编译 kernel 源码 → 产出 .so（host-wrapper）	50	│ ├─ TileLang 编译器编译 kernel 源码 → 产出 .so（host-wrapper）
51	│ └─ dlopen .so + dlsym("call") → 缓存函数指针（临时文件读取后立即 unlink）	51	│ └─ dlopen .so + dlsym("call") → 缓存函数指针（临时文件读取后立即 unlink）
52	└─ 返回 GRAPH_SUCCESS	52	└─ 返回 GRAPH_SUCCESS
	@@ -175,7 +175,7 @@ GE 原生构图测试程序：
175	\| 算子类型 \| `AddCustomOnline` \|	175	\| 算子类型 \| `AddCustomOnline` \|
176	\| 输入 \| `x` (float32), `y` (float32) \|	176	\| 输入 \| `x` (float32), `y` (float32) \|
177	\| 输出 \| `z` (float32) \|	177	\| 输出 \| `z` (float32) \|
178	-\| 输入 shape \| `[4096]` (固定) \|	178	+\| 输入 shape \| `[4096]`（固定） \|
179	\| 输出 shape \| `[4096]` \|	179	\| 输出 shape \| `[4096]` \|
180	\| 格式 \| ND \|	180	\| 格式 \| ND \|
181	\| kernel 名称 \| `main_kernel`（由 `call` 封装） \|	181	\| kernel 名称 \| `main_kernel`（由 `call` 封装） \|
	@@ -198,7 +198,7 @@ export ASCEND_CUSTOM_OPP_PATH="$(pwd)/output:$ASCEND_CUSTOM_OPP_PATH"
198		198
199	## 注意事项	199	## 注意事项
200		200
201	-- 同机有卡编译限定：TileLang-Ascend 当前通过 `torch.npu.get_device_name()` 做运行时平台检测，不支持离线指定目标架构。本样例仅适用于"编译机与目标机为同一 NPU"的场景。	201	+- 同机有卡编译限定：TileLang-Ascend 当前通过 `torch.npu.get_device_name()` 做运行时平台检测，不支持离线指定目标架构。本样例仅适用于“编译机与目标机为同一 NPU”的场景。
202	- kernel 源码 `.py` 安装在 OPP 包的 `op_graph/lib/<os>/<arch>/` 目录下，与 `libcust_opapi.so` 同目录，`Compile` 通过 `dladdr` 定位。	202	- kernel 源码 `.py` 安装在 OPP 包的 `op_graph/lib/<os>/<arch>/` 目录下，与 `libcust_opapi.so` 同目录，`Compile` 通过 `dladdr` 定位。
203	- 编译产出的 `.so` 使用 `mkstemps` 生成唯一临时文件，读取后立即 `unlink`，不会残留。	203	- 编译产出的 `.so` 使用 `mkstemps` 生成唯一临时文件，读取后立即 `unlink`，不会残留。
204	- `ge.graphRunMode=1` 确保走在线执行链路（PRIORITY_GRAPH 模式）。	204	- `ge.graphRunMode=1` 确保走在线执行链路（PRIORITY_GRAPH 模式）。

Mexamples/custom_op/tilelang_add_custom_online/README_en.md+121-28

	@@ -4,21 +4,21 @@
4		4
5	- Graph construction entry: GE native (Session API)	5	- Graph construction entry: GE native (Session API)
6	- Operator programming language: TileLang	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	7	+- Compilation method: GE compile phase invokes the 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`	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`)	9	- Scenario: Scenario B — online compilation + online execution (`CompilableOp` + `EagerExecuteOp` + `ShapeInferOp`)
10		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.	11	+This sample uses an element-wise Add operator to demonstrate how to compile TileLang kernel source online during GE's compile phase (`CompileGraph`) via the `CompilableOp` interface, rather than pre-compiling the `.so` and loading it. Contrast with the [tilelang_add_custom](../tilelang_add_custom/README_en.md) (eager mode) sample.
12		12
13	## Differences from Eager Sample	13	## Differences from Eager Sample
14		14
15	\| Dimension \| Eager (`tilelang_add_custom`) \| Online Compilation (this sample) \|	15	\| Dimension \| Eager (`tilelang_add_custom`) \| Online Compilation (this sample) \|
16	\|-----------\|-------------------------------\|----------------------------------\|	16	\|-----------\|-------------------------------\|----------------------------------\|
17	\| Interface combo \| `EagerExecuteOp` + `ShapeInferOp` \| `CompilableOp` + `EagerExecuteOp` + `ShapeInferOp` \|	17	\| Interface combo \| `EagerExecuteOp` + `ShapeInferOp` \| `CompilableOp` + `EagerExecuteOp` + `ShapeInferOp` \|
18	-\| Compilation timing \| Pre-compiled by `run.sh` \| Online during `CompileGraph` \|	18	+\| Compilation timing \| `run.sh` pre-compiles `.so` \| Online during `CompileGraph` \|
19	-\| Load timing \| Lazy `dlopen` at first `Execute` \| `dlopen` in `Compile`, cached for `Execute` \|	19	+\| Load timing \| First `Execute` call triggers lazy `dlopen` \| `dlopen` in `Compile` callback, `Execute` uses the cache directly \|
20	\| Compilation trigger \| Manual `python3 add_custom_kernel.py` \| GE `CustomGraphOptimizer` calls `Compile` \|	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 \|	21	+\| Shape caching \| None (fixed N=4096) \| Keyed by element count, supports multiple input sizes \|
22	\| Thread safety \| `std::once_flag` \| `std::mutex` (`Compile` may be called in parallel) \|	22	\| Thread safety \| `std::once_flag` \| `std::mutex` (`Compile` may be called in parallel) \|
23		23
24	## Directory Structure	24	## Directory Structure
	@@ -27,57 +27,76 @@ This sample demonstrates how to compile TileLang kernel source online during GE'
27	tilelang_add_custom_online/	27	tilelang_add_custom_online/
28	├── README.md	28	├── README.md
29	├── README_en.md	29	├── README_en.md
30	-├── CMakeLists.txt	30	+├── CMakeLists.txt # Build libcust_opapi.so + session_run + install .py source
31	-├── run.sh	31	+├── run.sh # One-click build and run (no kernel pre-compilation)
32	├── add_custom_kernel/	32	├── add_custom_kernel/
33	-│ └── add_custom_kernel.py	33	+│ └── add_custom_kernel.py # TileLang kernel source (accepts N and output path arguments)
34	├── ge/	34	├── ge/
35	-│ ├── add_custom.h	35	+│ ├── add_custom.h # REG_OP proto definition
36	-│ └── custom_op.cpp	36	+│ └── custom_op.cpp # CompilableOp + EagerExecuteOp + ShapeInferOp implementation
37	└── session_run/	37	└── session_run/
38	- └── main.cc	38	+ └── main.cc # GE native graph + CompileGraph (triggers online compilation) + execution + precision check
39	```	39	```
40		40
41	## Core Pipeline	41	## Core Pipeline
42		42
43	```text	43	```text
44	GE compile phase (CompileGraph):	44	GE compile phase (CompileGraph):
45	- CustomGraphOptimizer calls Compile(ctx)	45	+ CustomGraphOptimizer calls back Compile(ctx)
46	- ├─ Read input shape → build binary key	46	+ ├─ Read input element count → build binary key
47	├─ If key not cached:	47	├─ If key not cached:
48	│ ├─ Locate add_custom_kernel.py (in OPP package, same dir as libcust_opapi.so)	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>")	49	+ │ ├─ popen("python3 add_custom_kernel.py <N> <output.so>") (same-machine NPU compilation)
50	- │ ├─ TileLang compiler compiles kernel source → .so (host-wrapper)	50	+ │ ├─ TileLang compiler compiles kernel source → produces .so (host-wrapper)
51	- │ └─ dlopen .so + dlsym("call") → cache function pointer	51	+ │ └─ dlopen .so + dlsym("call") → cache function pointer (temp file unlinked immediately after reading)
52	└─ Return GRAPH_SUCCESS	52	└─ Return GRAPH_SUCCESS
53		53
54	GE execution phase (ExecuteGraphWithStreamAsync):	54	GE execution phase (ExecuteGraphWithStreamAsync):
55	- Execute(ctx) called	55	+ Execute(ctx) called back
56	- ├─ Read input shape → build binary key	56	+ ├─ Read input element count → build binary key
57	├─ Get cached call function pointer	57	├─ Get cached call function pointer
58	├─ Allocate output Tensor	58	├─ Allocate output Tensor
59	└─ call(x_ptr, y_ptr, z_ptr, stream) → NPU execution	59	└─ call(x_ptr, y_ptr, z_ptr, stream) → NPU execution
60	```	60	```
61		61
		62	+The `.so` compiled by TileLang-Ascend exports a function with the signature:
		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	+This function internally wraps the launch logic of `main_kernel<<<>>>` (including hardware scheduling address acquisition, tiling, etc.), so GE does not need to assemble args manually.
		69	+
62	## Prerequisites	70	## Prerequisites
63		71
64	### CANN	72	### CANN
65		73
66	- CANN environment properly installed and configured (`source ${ASCEND_HOME_PATH}/set_env.sh`)	74	- CANN environment properly installed and configured (`source ${ASCEND_HOME_PATH}/set_env.sh`)
		75	+- The environment provides ACL, GE, and Graph related headers and libraries
67		76
68	### TileLang-Ascend	77	### TileLang-Ascend
69		78
		79	+Install the TileLang main package and the TileLang-Ascend backend:
		80	+
70	```bash	81	```bash
71	pip install tilelang	82	pip install tilelang
72	# TileLang-Ascend backend: install from https://github.com/tile-ai/tilelang-ascend	83	# TileLang-Ascend backend: install from https://github.com/tile-ai/tilelang-ascend
73	```	84	```
74		85
75	-If installed from source, set:	86	+If TileLang-Ascend is installed from source (not via `pip install`), set the environment variable:
76		87
77	```bash	88	```bash
78	export TILELANG_ASCEND_HOME=/path/to/tilelang-ascend	89	export TILELANG_ASCEND_HOME=/path/to/tilelang-ascend
79	```	90	```
80		91
		92	+### Environment Variables
		93	+
		94	+\| Variable \| Required \| Description \|
		95	+\|----------\|----------\|-------------\|
		96	+\| `ASCEND_HOME_PATH` \| Yes \| CANN toolkit path \|
		97	+\| `TILELANG_ASCEND_HOME` \| No \| TileLang-Ascend source installation path (not needed if installed via pip) \|
		98	+\| `ASCEND_CUSTOM_OPP_PATH` \| Auto \| Set automatically by `run.sh` \|
		99	+
81	## Quick Start	100	## Quick Start
82		101
83	```bash	102	```bash
	@@ -85,14 +104,70 @@ source ${ASCEND_HOME_PATH}/set_env.sh
85	bash run.sh	104	bash run.sh
86	```	105	```
87		106
88	-`run.sh` executes 3 steps:	107	+`run.sh` executes 3 steps in sequence:
89		108
90	-1. Build `libcust_opapi.so` and `tilelang_online_session_run`, install `add_custom_kernel.py` to OPP package	109	+1. Build `libcust_opapi.so` and `tilelang_online_session_run`, install `add_custom_kernel.py` into the OPP package
91	-2. Verify kernel source is in OPP package	110	+2. Verify the kernel source is in the OPP package
92	-3. Run test program (`CompileGraph` triggers TileLang online compilation, then executes and verifies)	111	+3. Run the test program (`CompileGraph` triggers TileLang online compilation, then executes and verifies precision)
93		112
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.	113	> 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		114
		115	+Expected terminal output on success:
		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	+## Key Files
		130	+
		131	+### `ge/custom_op.cpp`
		132	+
		133	+GE deliverable implementing `CompilableOp` + `EagerExecuteOp` + `ShapeInferOp`:
		134	+
		135	+- Compile:
		136	+ 1. Read the input shape size from `ctx->GetInputTensor(0)` and build a binary key
		137	+ 2. Lock and check the cache; if the key already exists, return directly (multiple shapes supported)
		138	+ 3. Locate `add_custom_kernel.py` via `dladdr` on the directory of `libcust_opapi.so`
		139	+ 4. Call `python3 add_custom_kernel.py <N> <output.so>` via `popen` to compile the TileLang kernel
		140	+ 5. `dlopen` the compiled `.so`, obtain the function pointer via `dlsym("call")`, and cache it in `kernel_entries_`
		141	+- Execute:
		142	+ 1. Read the input shape size and build the key
		143	+ 2. Get the function pointer cached during `Compile` from `kernel_entries_`
		144	+ 3. Allocate the output Tensor and call `call(x_ptr, y_ptr, z_ptr, stream)`
		145	+- InferShape / InferDataType: output shape and dtype are the same as the input
		146	+- Uses `std::mutex` for thread safety (`CustomGraphOptimizer` may call `Compile` in parallel)
		147	+
		148	+### `ge/add_custom.h`
		149	+
		150	+`REG_OP(AddCustomOnline)` declares the operator's input/output specification, used by GE native graph construction to create nodes.
		151	+
		152	+### `add_custom_kernel/add_custom_kernel.py`
		153	+
		154	+TileLang kernel source, accepting command-line arguments:
		155	+
		156	+- 1st argument: `N` (total element count, default 4096, must be a multiple of BLOCK_SIZE=1024)
		157	+- 2nd argument: `output_path` (path of the produced `.so`)
		158	+
		159	+### `session_run/main.cc`
		160	+
		161	+GE native graph construction test program:
		162	+
		163	+1. `GEInitialize` + create a `Session`
		164	+2. Build the `Data → AddCustomOnline` computation graph
		165	+3. `AddGraph` → `CompileGraph` (triggers `CompilableOp::Compile` → TileLang online compilation)
		166	+4. `LoadGraph`
		167	+5. Allocate device memory, H2D copy of input data
		168	+6. Execute with `ExecuteGraphWithStreamAsync`
		169	+7. D2H copy of the output, element-wise precision check (including NaN check)
		170	+
96	## Operator Specification	171	## Operator Specification
97		172
98	\| Item \| Value \|	173	\| Item \| Value \|
	@@ -101,15 +176,33 @@ bash run.sh
101	\| Inputs \| `x` (float32), `y` (float32) \|	176	\| Inputs \| `x` (float32), `y` (float32) \|
102	\| Output \| `z` (float32) \|	177	\| Output \| `z` (float32) \|
103	\| Input shape \| `[4096]` (fixed) \|	178	\| Input shape \| `[4096]` (fixed) \|
		179	+\| Output shape \| `[4096]` \|
104	\| Format \| ND \|	180	\| Format \| ND \|
105	\| Kernel name \| `main_kernel` (wrapped by `call`) \|	181	\| Kernel name \| `main_kernel` (wrapped by `call`) \|
106	\| BLOCK_SIZE \| 1024 \|	182	\| BLOCK_SIZE \| 1024 \|
107		183
		184	+## Step-by-Step Run
		185	+
		186	+```bash
		187	+# 1. Build (including installing add_custom_kernel.py into the OPP package)
		188	+cmake -S . -B build -DCMAKE_BUILD_TYPE=Release
		189	+cmake --build build -j$(nproc)
		190	+cmake --install build
		191	+
		192	+# 2. Set environment variables
		193	+export ASCEND_CUSTOM_OPP_PATH="$(pwd)/output:$ASCEND_CUSTOM_OPP_PATH"
		194	+
		195	+# 3. Run (CompileGraph triggers online compilation)
		196	+./build/tilelang_online_session_run
		197	+```
		198	+
108	## Notes	199	## Notes
109		200
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.	201	+- Same-machine NPU compilation required: TileLang-Ascend currently uses `torch.npu.get_device_name()` for runtime platform detection and does not support specifying the target architecture offline. This sample only applies to scenarios where the compilation machine and the target machine have the same NPU.
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`.	202	+- The 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.	203	+- The compiled `.so` uses `mkstemps` for a unique temp file and calls `unlink` immediately after reading, leaving no residue.
113	-- `ge.graphRunMode=1` ensures online execution (PRIORITY_GRAPH mode).	204	+- `ge.graphRunMode=1` ensures the online execution path (PRIORITY_GRAPH mode).
114	- `CompileGraph` must be called before `ExecuteGraphWithStreamAsync`, otherwise `Execute` cannot find the compiled kernel.	205	- `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.	206	+- Only float32 is supported. To support more data types, adjust the `REG_OP` `DATATYPE` constraint and the TileLang kernel dtype parameter.
		207	+- TileLang-Ascend platform detection is based on `torch.npu.get_device_name()`. Ascend910 maps to the A2 platform.
		208	+- Online compilation requires Python + TileLang in the runtime environment, suitable for the development phase; for production deployment, consider the eager sample's pre-compilation approach.