Pull Request

变更摘要

本 PR 在 examples/custom_op/ 下新增三个 TileLang 自定义算子样例，分别演示如何用 TileLang 开发 Ascend 自定义算子并通过 GE 完成入图、在线编译和离线 OM 模型下沉，同时更新了 examples/custom_op/README.md 与 README_en.md 的样例索引。三个样例共用同一份 TileLang element-wise Add kernel 源码：由 tilelang.jit 装饰的 vec_add 编译产出导出 call(uint8_t* A, uint8_t* B, uint8_t* C, aclrtStream stream) 签名的 .so，GE 侧通过 dlopen/dlsym("call") 加载并调用。

主要改动

• TileLang kernel 编译脚本（三个样例均含 add_custom_kernel/add_custom_kernel.py）：使用 T.Kernel、T.alloc_ub、T.tile.add 实现 element-wise Add kernel，经 tilelang.jit 编译后把生成的 .so 拷贝到指定路径；online/offline 版本支持 [N] [output_path] 命令行参数，便于运行时按 shape 动态编译。

• 入图样例 tilelang_add_custom：ge/add_custom.h 通过 REG_OP(AddCustom) 注册算子（INPUT x/y、OUTPUT z、DT_FLOAT）；ge/custom_op.cpp 中 AddCustom 继承 EagerExecuteOp/ShapeInferOp，按 ASCEND_CUSTOM_OPP_PATH 定位 add_kernel.so 并 dlopen/dlsym("call")，Execute 校验输入元素数（4096）后调用 kernel；session_run/main.cc 构建 Data→AddCustom 图，经 AddGraph/CompileGraph/LoadGraph/ExecuteGraphWithStreamAsync 执行并做精度校验。

• 在线编译样例 tilelang_add_custom_online：AddCustomOnline 继承 CompilableOp/EagerExecuteOp/ShapeInferOp，Compile 通过 fork + execvp 子进程调用 python3 add_custom_kernel.py <n> <so_path> 完成运行时编译，以 shape size 为 key 缓存并 dlopen 加载；Execute 按 key 查找已编译 kernel 后调用。

• 离线下沉样例 tilelang_add_custom_offline：AddCustomOffline 额外继承 PortableOp，Serialize 将编译产出的 kernel .so 二进制（含 magic/version/条目数）写入模型缓冲，Deserialize 校验 kMaxEntryCount/kMaxSoSize 等上限后经 memfd_create 创建内存文件再 dlopen；配套 graph_build 导出 AIR 文件，run.sh 调用 ATC 将 AIR 编译为 OM（默认 soc_version 为 Ascend910_9362），model_exec 通过 aclmdlLoadFromFile/aclmdlExecute 加载执行并校验结果。

代码审查

✅ 未发现问题

/lgtm

/approve