新增算子开发指南-pyasc-AtomGit

已合并

新增算子开发指南 #87

li-yuanjie-da创建于 3月17日从li-yuanjie-da/pyasc_1： master 合入到 cann/pyasc： master

新增算子开发指南 #87

已合并

li-yuanjie-da创建于 3月17日

从li-yuanjie-da/pyasc_1 master合入到cann/pyasc master

共 5 个文件变更+343-2

MREADME.md+1-0

	@@ -51,6 +51,7 @@ pyasc编程接口与Ascend C类库接口一一对应，旨在提供与Ascend C
51	- [Ascend C Python编程接口开发指南](https://gitcode.com/cann/pyasc/blob/master/docs/developer_guide.md)：介绍如何开发Ascend C API的Python编程接口。	51	- [Ascend C Python编程接口开发指南](https://gitcode.com/cann/pyasc/blob/master/docs/developer_guide.md)：介绍如何开发Ascend C API的Python编程接口。
52	- [pyasc的python语法支持情况说明](https://gitcode.com/cann/pyasc/blob/master/docs/python_syntax_support.md)：介绍pyasc项目支持和不支持的python语法。	52	- [pyasc的python语法支持情况说明](https://gitcode.com/cann/pyasc/blob/master/docs/python_syntax_support.md)：介绍pyasc项目支持和不支持的python语法。
53	- [API文档自动生成工具使用指南](https://gitcode.com/cann/pyasc/blob/master/docs/API_docstring_generation_tool_guide.md)：介绍本项目接口文档的生成方法。	53	- [API文档自动生成工具使用指南](https://gitcode.com/cann/pyasc/blob/master/docs/API_docstring_generation_tool_guide.md)：介绍本项目接口文档的生成方法。
		54	+- [PyAsc算子开发指南](https://gitcode.com/cann/pyasc/blob/master/docs/pyasc_op_develop_guide.md)：介绍基于pyasc的Ascend C算子开发基本流程。
54		55
55	## 👥 合作贡献者	56	## 👥 合作贡献者
56		57

Mdocs/README.md+1-0

	@@ -34,6 +34,7 @@
34	\| [Ascend C Python算子调试调优指南](op_debug_prof.md) \| 介绍常见的算子调试和调优方法。 \|	34	\| [Ascend C Python算子调试调优指南](op_debug_prof.md) \| 介绍常见的算子调试和调优方法。 \|
35	\| [pyasc的python语法支持情况说明](python_syntax_support.md) \| 介绍pyasc项目支持和不支持的python语法。 \|	35	\| [pyasc的python语法支持情况说明](python_syntax_support.md) \| 介绍pyasc项目支持和不支持的python语法。 \|
36	\| [API文档自动生成工具使用指南](API_docstring_generation_tool_guide.md) \| 介绍自动生成API文档的工具的使用方法。 \|	36	\| [API文档自动生成工具使用指南](API_docstring_generation_tool_guide.md) \| 介绍自动生成API文档的工具的使用方法。 \|
		37	+\| [PyAsc算子开发指南](pyasc_op_develop_guide.md) \| 介绍基于pyasc的Ascend C算子开发基本流程。 \|
37		38
38	## 附录	39	## 附录
39		40

Adocs/figures/pyasc_develop_guide.png+0-0

Adocs/pyasc_op_develop_guide.md+339-0

@@ -0,0 +1,339 @@
	1	+# PyAsc算子开发指南
	A A ai_xin3月19日首页readme、docs readme建议有个说明和跳转到本文档的链接
	2	+
	3	+本文档从一个简单的算子开发样例出发，带您体验基于pyasc的Ascend C算子开发基本流程。
	4	+
	5	+> 注意：本教程基于 [pyasc](https://gitcode.com/cann/pyasc) 项目，使用 Python 原生语法编写 Ascend C 算子，与 `.asc` 文件开发方式不同。
	6	+
	7	+在正式的开发之前，需要先完成环境准备工作，开发 pyasc 算子的基本流程如下：
	8	+
	9	+![开发流程](figures/pyasc_develop_guide.png)
	10	+
	11	+## 环境准备
	12	+
	13	+- pyasc 安装
	14	+
	15	+ pyasc 支持通过 pip 快速安装和基于源码编译安装两种方式。具体请参考 [pyasc 快速入门-编译环境准备](quick_start.md#buildenv)。
	16	+
	17	+- CANN 软件安装
	18	+
	19	+ 开发算子前需要安装 CANN 软件。安装 CANN 软件后，需要设置环境变量。具体请参考 [pyasc 快速入门-运行环境准备](quick_start.md#runtimeenv)。
	20	+
	21	+## 算子分析
	22	+
	23	+主要分析算子的数学表达式、输入输出的数量、Shape 范围以及计算逻辑的实现，明确需要调用的 pyasc 接口。下文以 Add 算子为例，介绍具体的分析过程。
	24	+
	25	+1. 明确算子的数学表达式及计算逻辑
	26	+
	27	+ Add 算子的数学表达式为：
	28	+ ```
	29	+ z = x + y
	30	+ ```
	31	+
	32	+ 计算逻辑是：从外部存储 Global Memory 搬运数据至内部存储 Local Memory，然后使用 pyasc 计算接口完成两个输入参数相加，得到最终结果，再搬运到 Global Memory 上。
	33	+
	34	+2. 明确输入和输出
	35	+
	36	+ - Add 算子有两个输入：x 与 y，输出为 z
	37	+ - 本样例中算子输入支持的数据类型为 float，算子输出的数据类型与输入数据类型相同
	38	+ - 算子输入支持的 shape 为（8，2048），输出 shape 与输入 shape 相同
	39	+ - 算子输入支持的 format 为：ND
	40	+
	41	+3. 确定核函数名称和参数
	42	+
	43	+ - 本样例中核函数命名为 `vadd_kernel`
	44	+ - 根据对算子输入输出的分析，确定核函数有 3 个参数 x，y，z；x，y 为输入参数，z 为输出参数
	45	+
	46	+4. 确定算子实现所需接口
	K K kong08083月17日这里可以加一个描述，相关的算子调用接口具体可参考Python编程接口列表
	47	+
	48	+ - 实现涉及外部存储和内部存储间的数据搬运，使用 [`asc.data_copy`](./python-api/language/generated/asc.language.basic.data_copy.md) 接口来实现数据搬移
	49	+ - 本样例只涉及矢量计算的加法操作，使用 [`asc.add`](./python-api/language/generated/asc.language.basic.add.md) 接口实现 x+y
	50	+ - 计算中使用到的 Tensor 数据结构，使用 [`asc.GlobalTensor`](./python-api/language/core.md)、[`asc.LocalTensor`](./python-api/language/core.md) 进行管理
	51	+ - 并行流水任务之间使用 [`asc.set_flag`](./python-api/language/generated/asc.language.basic.set_flag.md)/[`asc.wait_flag`](./python-api/language/generated/asc.language.basic.wait_flag.md) 接口完成同步
	52	+
	53	+通过以上分析，得到 pyasc Add 算子的设计规格如下：
	54	+
	55	+ <table>
	56	+ <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="5" align="center">Add</td></tr>
	57	+ </tr>
	58	+ <tr><td rowspan="3" align="center">算子输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>
	59	+ <tr><td align="center">x</td><td align="center">(8, 2048)</td><td align="center">float</td><td align="center">ND</td></tr>
	60	+ <tr><td align="center">y</td><td align="center">(8, 2048)</td><td align="center">float</td><td align="center">ND</td></tr>
	61	+ </tr>
	62	+ </tr>
	63	+ <tr><td rowspan="1" align="center">算子输出</td><td align="center">z</td><td align="center">(8, 2048)</td><td align="center">float</td><td align="center">ND</td></tr>
	64	+ </tr>
	65	+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="5" align="center">vadd_kernel</td></tr>
	66	+ <tr><td rowspan="4" align="center">使用的主要接口</td><td colspan="5" align="center">asc.data_copy：数据搬运接口</td></tr>
	67	+ <tr><td colspan="5" align="center">asc.add：矢量基础算术接口</td></tr>
	68	+ <tr><td colspan="5" align="center">asc.GlobalTensor/LocalTensor：内存管理接口</td></tr>
	69	+ <tr><td colspan="5" align="center">asc.set_flag/wait_flag：同步接口</td></tr>
	70	+ <tr><td rowspan="1" align="center">算子实现文件名称</td><td colspan="5" align="center">add.py</td></tr>
	71	+ </table>
	72	+
	73	+---
	74	+
	75	+## 核函数开发
	76	+
	77	+完成环境准备和初步的算子分析后，即可开始 pyasc 核函数的开发。
	78	+
	79	+本样例中使用多核并行计算，即把数据进行分片，分配到多个核上进行处理。pyasc 核函数是在一个核上的处理函数，所以只处理部分数据。分配方案是：假设共启用 8 个核，数据整体长度为 8 * 2048 个元素，平均分配到 8 个核上运行，每个核上处理的数据大小为 2048 个元素。对于单核上的处理数据，也可以进行数据切块，实现对数据的流水并行处理。
	80	+
	81	+1. 定义核函数参数
	82	+
	83	+ 本样例使用以下参数控制数据切分：
	84	+ - `USE_CORE_NUM = 8`：启用 8 个核
	85	+ - `TILE_NUM = 8`：每个核上数据分块个数
	86	+ - `BUFFER_NUM = 2`：双缓冲
	87	+
	88	+2. 核函数定义与实现
	89	+
	90	+ 使用 `@asc.jit` 装饰器定义核函数，并在核函数中实现算子逻辑：
	91	+
	92	+ ```python
	93	+ import asc
	94	+ import asc.lib.runtime as rt
	95	+
	96	+ USE_CORE_NUM = 8
	97	+ BUFFER_NUM = 2
	98	+ TILE_NUM = 8
	99	+
	100	+ @asc.jit
	101	+ def vadd_kernel(x: asc.GlobalAddress, y: asc.GlobalAddress, z: asc.GlobalAddress, block_length: int):
	102	+ # 获取当前核的索引，计算数据偏移
	103	+ offset = asc.get_block_idx() * block_length
	104	+
	105	+ # 创建 GlobalTensor 管理全局内存地址
	106	+ x_gm = asc.GlobalTensor()
	107	+ y_gm = asc.GlobalTensor()
	108	+ z_gm = asc.GlobalTensor()
	109	+
	110	+ # 设置 Global Memory 起始地址和长度
	111	+ x_gm.set_global_buffer(x + offset, block_length)
	112	+ y_gm.set_global_buffer(y + offset, block_length)
	113	+ z_gm.set_global_buffer(z + offset, block_length)
	114	+
	115	+ # 计算每个 tile 的长度（考虑双缓冲）
	116	+ tile_length = block_length // TILE_NUM // BUFFER_NUM
	117	+
	118	+ # 获取数据类型信息
	119	+ data_type = x.dtype
	120	+ buffer_size = tile_length * BUFFER_NUM * data_type.sizeof()
	121	+
	122	+ # 创建 LocalTensor（基于指定的逻辑位置/地址/长度）
	123	+ # x_local 和 y_local 放在 VECIN 位置
	124	+ x_local = asc.LocalTensor(data_type, asc.TPosition.VECIN, 0, tile_length * BUFFER_NUM)
	125	+ y_local = asc.LocalTensor(data_type, asc.TPosition.VECIN, buffer_size, tile_length * BUFFER_NUM)
	126	+ # z_local 放在 VECOUT 位置
	127	+ z_local = asc.LocalTensor(data_type, asc.TPosition.VECOUT, buffer_size + buffer_size, tile_length * BUFFER_NUM)
	128	+
	129	+ # 流水循环处理（双缓冲需要循环次数翻倍）
	130	+ for i in range(TILE_NUM * BUFFER_NUM):
	131	+ buf_id = i % BUFFER_NUM
	132	+
	133	+ # Step 1: 搬入 - 从 Global Memory 拷贝数据到 Local Memory
	134	+ asc.data_copy(x_local[buf_id * tile_length:], x_gm[i * tile_length:], tile_length)
	135	+ asc.data_copy(y_local[buf_id * tile_length:], y_gm[i * tile_length:], tile_length)
	136	+
	137	+ # 同步：等待 MTE2_V 事件，确保数据搬入完成
	138	+ asc.set_flag(asc.HardEvent.MTE2_V, buf_id)
	139	+ asc.wait_flag(asc.HardEvent.MTE2_V, buf_id)
	140	+
	141	+ # Step 2: 计算 - 执行矢量加法
	142	+ asc.add(z_local[buf_id * tile_length:], x_local[buf_id * tile_length:],
	143	+ y_local[buf_id * tile_length:], tile_length)
	144	+
	145	+ # 同步：等待 V_MTE3 事件，确保计算完成
	146	+ asc.set_flag(asc.HardEvent.V_MTE3, buf_id)
	147	+ asc.wait_flag(asc.HardEvent.V_MTE3, buf_id)
	148	+
	149	+ # Step 3: 搬出 - 从 Local Memory 拷贝数据到 Global Memory
	150	+ asc.data_copy(z_gm[i * tile_length:], z_local[buf_id * tile_length:], tile_length)
	151	+
	152	+ # 同步：等待 MTE3_MTE2 事件，确保数据搬出完成
	153	+ asc.set_flag(asc.HardEvent.MTE3_MTE2, buf_id)
	154	+ asc.wait_flag(asc.HardEvent.MTE3_MTE2, buf_id)
	155	+ ```
	156	+
	157	+ 内部函数的调用关系示意图：
	158	+
	159	+ ```
	160	+ vadd_kernel
	161	+ ├── offset = get_block_idx() * block_length
	162	+ ├── GlobalTensor 设置
	163	+ │ ├── x_gm.set_global_buffer()
	164	+ │ ├── y_gm.set_global_buffer()
	165	+ │ └── z_gm.set_global_buffer()
	166	+ ├── LocalTensor 创建
	167	+ └── for i in range(TILE_NUM * BUFFER_NUM):
	168	+ ├── CopyIn: data_copy (x_local, y_local <- x_gm, y_gm)
	169	+ ├── Compute: add (z_local <- x_local + y_local)
	170	+ └── CopyOut: data_copy (z_gm <- z_local)
	171	+ ```
	172	+
	173	+3. Launch 函数实现
	174	+
	175	+ ```python
	176	+ def vadd_launch(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
	177	+ z = torch.zeros_like(x)
	178	+
	179	+ total_length = z.numel()
	180	+ block_length = total_length // USE_CORE_NUM
	181	+
	182	+ vadd_kernel[USE_CORE_NUM, rt.current_stream()](x, y, z, block_length)
	183	+ return z
	184	+ ```
	185	+
	186	+ - `vadd_kernel[USE_CORE_NUM, rt.current_stream()]`：使用内核调用符指定核数和流
	187	+ - `(x, y, z, block_length)`：传递参数
	188	+
	189	+---
	190	+
	191	+## 核函数运行验证
	192	+
	193	+完成核函数开发后，即可编写完整的核函数调用程序，执行计算过程。
	A A ai_xin3月19日 vadd_custom的两个入参能否说明一下含义作用来源
	194	+
	195	+1. 完整的算子验证程序
	196	+
	197	+ ```python
	198	+ import logging
	199	+ import argparse
	200	+ import torch
	201	+ try:
	202	+ import torch_npu
	203	+ except ModuleNotFoundError:
	204	+ pass
	205	+
	206	+ import asc
	207	+ import asc.runtime.config as config
	208	+ import asc.lib.runtime as rt
	209	+
	210	+ USE_CORE_NUM = 8
	211	+ BUFFER_NUM = 2
	212	+ TILE_NUM = 8
	213	+
	214	+ logging.basicConfig(level=logging.INFO)
	215	+
	216	+ @asc.jit
	217	+ def vadd_kernel(x: asc.GlobalAddress, y: asc.GlobalAddress, z: asc.GlobalAddress, block_length: int):
	218	+ offset = asc.get_block_idx() * block_length
	219	+ x_gm = asc.GlobalTensor()
	220	+ y_gm = asc.GlobalTensor()
	221	+ z_gm = asc.GlobalTensor()
	222	+ x_gm.set_global_buffer(x + offset, block_length)
	223	+ y_gm.set_global_buffer(y + offset, block_length)
	224	+ z_gm.set_global_buffer(z + offset, block_length)
	225	+
	226	+ tile_length = block_length // TILE_NUM // BUFFER_NUM
	227	+
	228	+ data_type = x.dtype
	229	+ buffer_size = tile_length * BUFFER_NUM * data_type.sizeof()
	230	+
	231	+ x_local = asc.LocalTensor(data_type, asc.TPosition.VECIN, 0, tile_length * BUFFER_NUM)
	232	+ y_local = asc.LocalTensor(data_type, asc.TPosition.VECIN, buffer_size, tile_length * BUFFER_NUM)
	233	+ z_local = asc.LocalTensor(data_type, asc.TPosition.VECOUT, buffer_size + buffer_size, tile_length * BUFFER_NUM)
	234	+
	235	+ for i in range(TILE_NUM * BUFFER_NUM):
	236	+ buf_id = i % BUFFER_NUM
	237	+
	238	+ asc.data_copy(x_local[buf_id * tile_length:], x_gm[i * tile_length:], tile_length)
	239	+ asc.data_copy(y_local[buf_id * tile_length:], y_gm[i * tile_length:], tile_length)
	240	+
	241	+ asc.set_flag(asc.HardEvent.MTE2_V, buf_id)
	242	+ asc.wait_flag(asc.HardEvent.MTE2_V, buf_id)
	243	+
	244	+ asc.add(z_local[buf_id * tile_length:], x_local[buf_id * tile_length:],
	245	+ y_local[buf_id * tile_length:], tile_length)
	246	+
	247	+ asc.set_flag(asc.HardEvent.V_MTE3, buf_id)
	248	+ asc.wait_flag(asc.HardEvent.V_MTE3, buf_id)
	249	+
	250	+ asc.data_copy(z_gm[i * tile_length:], z_local[buf_id * tile_length:], tile_length)
	251	+
	252	+ asc.set_flag(asc.HardEvent.MTE3_MTE2, buf_id)
	253	+ asc.wait_flag(asc.HardEvent.MTE3_MTE2, buf_id)
	254	+
	255	+ def vadd_launch(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
	256	+ z = torch.zeros_like(x)
	257	+ total_length = z.numel()
	258	+ block_length = total_length // USE_CORE_NUM
	259	+ vadd_kernel[USE_CORE_NUM, rt.current_stream()](x, y, z, block_length)
	260	+ return z
	261	+
	262	+ # Backend 对应执行脚本时传入的 [RUN_MODE]
	263	+ # Platform 对应执行脚本时传入的 [SOC_VERSION]
	264	+ def vadd_custom(backend: config.Backend, platform: config.Platform):
	265	+ config.set_platform(backend, platform)
	266	+ device = "npu" if config.Backend(backend) == config.Backend.NPU else "cpu"
	267	+ size = 8 * 2048
	268	+ x = torch.rand(size, dtype=torch.float32, device=device)
	269	+ y = torch.rand(size, dtype=torch.float32, device=device)
	270	+ z = vadd_launch(x, y)
	271	+ assert torch.allclose(z, x + y)
	272	+
	273	+ if __name__ == "__main__":
	274	+ parser = argparse.ArgumentParser()
	275	+ parser.add_argument("-r", type=str, default="Model", help="backend to run")
	276	+ parser.add_argument("-v", type=str, default=None, help="platform to run")
	277	+ args = parser.parse_args()
	278	+ backend = args.r
	279	+ platform = args.v
	280	+ if backend not in config.Backend.__members__:
	281	+ raise ValueError("Unsupported Backend! Supported: ['Model', 'NPU']")
	282	+ backend = config.Backend(backend)
	283	+ if platform is not None:
	284	+ platform_values = [platform.value for platform in config.Platform]
	285	+ if platform not in platform_values:
	286	+ raise ValueError(f"Unsupported Platform! Supported: {platform_values}")
	287	+ platform = config.Platform(platform)
	288	+ logging.info("[INFO] start process sample add.")
	289	+ vadd_custom(backend, platform)
	290	+ logging.info("[INFO] Sample add run success.")
	291	+ ```
	292	+
	293	+2. 编译和运行
	294	+
	295	+ 运行时使用以下命令：
	296	+
	297	+ ```bash
	298	+ python3 add.py -r [RUN_MODE] -v [SOC_VERSION]
	299	+ ```
	300	+
	301	+ 其中：
	302	+ - `RUN_MODE`：编译执行方式，可选择 `Model`（仿真）或 `NPU`（上板）。
	303	+ - `SOC_VERSION`：昇腾 AI 处理器型号，如果无法确定具体的[SOC_VERSION]，则在安装昇腾AI处理器的服务器执行npu-smi info命令进行查询，在查询到的“Name”前增加Ascend信息，例如“Name”对应取值为xxxyy，实际配置的[SOC_VERSION]值为Ascendxxxyy。
	304	+
	305	+ 示例：
	306	+ ```bash
	307	+ # 仿真器模式运行
	308	+ python3 add.py -r Model -v Ascend910B1
	309	+
	310	+ # NPU 上板模式运行
	311	+ python3 add.py -r NPU -v Ascend910B1
	312	+ ```
	313	+
	314	+ 用例执行完成，打屏信息出现 `Sample add run success.`，说明样例执行成功。
	315	+
	316	+---
	317	+
	318	+## pyasc 与 Ascend C 算子开发接口/语法特性的对比
	319	+
	320	+\| 特性 \| Ascend C (.asc) \| pyasc (Python) \|
	321	+\|------\|-----------------\|----------------\|
	322	+\| 编程语言 \| C++ 扩展语法 \| 原生 Python \|
	323	+\| 核函数定义 \| `__global__ __aicore__` \| `@asc.jit` 装饰器 \|
	324	+\| GlobalTensor \| `AscendC::GlobalTensor<T>` \| `asc.GlobalTensor()` \|
	325	+\| LocalTensor \| `AscendC::LocalTensor<T>` \| `asc.LocalTensor()` \|
	326	+\| 数据搬运 \| `AscendC::DataCopy()` \| `asc.data_copy()` \|
	327	+\| 矢量计算 \| `AscendC::Add()` \| `asc.add()` \|
	328	+\| 同步事件 \| `AscendC::SetFlag()/WaitFlag()` \| `asc.set_flag()/wait_flag()` \|
	329	+\| 核函数调用 \| `add_custom<<<...>>>()` \| `vadd_kernel[num_blocks, stream]()` \|
	330	+
	331	+---
	332	+
	333	+## 接下来的引导
	334	+
	335	+- 如果您想了解更多 pyasc 算子示例，可以参考 [tutorials](../python/tutorials) 目录下的样例。
	336	+
	337	+- 如果您想深入了解 pyasc 的 API 接口，请参考 [API 文档](./python-api/index.md)。
	338	+
	339	+- 如果您想了解 pyasc 的构建和调试方法，请参考 [快速入门](quick_start.md)。

Mdocs/quick_start.md+2-2

	@@ -1,6 +1,6 @@
1	# 构建	1	# 构建
2	## 环境准备<a name="envready"></a>	2	## 环境准备<a name="envready"></a>
3	-### 编译环境准备	3	+### 编译环境准备<a name="buildenv"></a>
4	pyasc支持通过pip快速安装和基于源码编译安装两种方式。	4	pyasc支持通过pip快速安装和基于源码编译安装两种方式。
5		5
6	#### 快速安装	6	#### 快速安装
	@@ -160,7 +160,7 @@ pyasc支持通过pip快速安装和基于源码编译安装两种方式。
160	python3 -m pip install -e .	160	python3 -m pip install -e .
161	```	161	```
162		162
163	-### 运行环境准备	163	+### 运行环境准备<a name="runtimeenv"></a>
164		164
165	使用基于源码安装时，建议安装社区版<a href="https://www.hiascend.com/developer/download/community/result?module=cann&cann=8.5.0.alpha001">8.5.0.alpha001</a>及以上版本。	165	使用基于源码安装时，建议安装社区版<a href="https://www.hiascend.com/developer/download/community/result?module=cann&cann=8.5.0.alpha001">8.5.0.alpha001</a>及以上版本。
166		166