[feat]add cpp_extension OpCommand guide-op-plugin-AtomGit

已合并

[feat]add cpp_extension OpCommand guide #5744

梁松伟创建于 14 天前从liangsongwei/op-plugin： master 合入到 Ascend/op-plugin： master

[feat]add cpp_extension OpCommand guide #5744

已合并

梁松伟创建于 14 天前

从liangsongwei/op-plugin master合入到Ascend/op-plugin master

共 4 个文件变更+214-69

Mexamples/cpp_extension/README.md+72-40

	@@ -1,12 +1,14 @@
1	-# 适配开发及调用	1	+# 自定义算子 C++ 扩展开发示例
2		2
3	-## 目录结构介绍	3	+本示例演示如何使用 Ascend C 实现自定义算子 Kernel，并通过 C++ 扩展（`cpp_extension`）方式集成到 PyTorch，最终在 Python 侧调用。涵盖工程结构、编译、安装、测试全流程。
		4	+
		5	+## 目录结构
4		6
5	```text	7	```text
6	├── examples	8	├── examples
7	\| ├── cpp_extension	9	\| ├── cpp_extension
8	\| │ ├── csrc	10	\| │ ├── csrc
9	-\| │ │ ├── add_custom.asc # Add算子实现	11	+\| │ │ ├── add_custom.asc # Add算子实现（含 5 种正确启动方式对比）
10	\| │ │ ├── trig_inplace_custom.asc # 原地三角函数算子实现	12	\| │ │ ├── trig_inplace_custom.asc # 原地三角函数算子实现
11	\| │ │ └── pybind11.asc # pybind绑定自定义算子	13	\| │ │ └── pybind11.asc # pybind绑定自定义算子
12	\| \| ├── op_extension/	14	\| \| ├── op_extension/
	@@ -21,45 +23,61 @@
21		23
22	## 新增自定义算子	24	## 新增自定义算子
23		25
24	-本示例以Add算子为例，展示了如何在 PyTorch 中使用Ascend C扩展自定义算子Kernel实现，并通过Python的接口调用实现的算子。	26	+### 1. Kernel 实现
25		27
26	-### kernel实现	28	+在 `./csrc/` 下创建 `.asc` 文件，基于 Ascend C 实现算子 Kernel。Ascend C 开发参考[昇腾社区文档](https://www.hiascend.com/ascend-c)。
27		29
28	- 本小节主要介绍如何实现Kernel算子，本样例基于Ascend C进行开发算子。如何使用Ascend C实现算子kernel，可以参考昇腾社区文档[昇腾Ascend C](https://www.hiascend.com/ascend-c)。	30	+以 `add_custom.asc` 为例，文件包含三部分：
29		31
30	- 在`./csrc/`目录下创建一个名为`add_custom.asc`的文件，这是我们自定义加法Kernel的实现文件。样例中实现了一个`run_ascendc_add`的核函数。	32	+\| 组成 \| 说明 \|
		33	+\|------\|------\|
		34	+\| `KernelAdd` 类 \| 算子设备侧实现，包含 `Init` / `Process` / `CopyIn` / `Compute` / `CopyOut` \|
		35	+\| `add_custom` 核函数 \| `__global__ __vector__` 修饰的设备入口 \|
		36	+\| `ascendc_add1/2/3/4/5` \| 5 种正确的 Host 侧启动方式，对比不同 stream 获取与队列管理策略 \|
31		37
32	-### 将Kernel实现与Python接口集成	38	+5 种启动方式对比：
33		39
34	- 本小节主要介绍如何封装实现的kernel算子，以及绑定为Python的接口。	40	+\| 方式 \| 函数 \| stream 获取 \| 队列管理 \| 适用场景 \|
35	- 代码实现在./csrc/ 目录下。	41	+\|------\|------\|------------\|---------\|---------\|
		42	+\| 1 \| `ascendc_add1` \| `NPUStream` 对象 \| `<<<>>>` 内部清 queue \| 简单同步场景 \|
		43	+\| 2 \| `ascendc_add2` \| `stream(true)` \| 清 queue 后直接启动 \| 等价方式1 \|
		44	+\| 3 \| `ascendc_add3` \| `stream(false)` \| OpCommand 入 queue \| 推荐，保留流水线性能 \|
		45	+\| 4 \| `ascendc_add4` \| `stream(true)` \| 清 queue + OpCommand 入 queue \| 语义直观 \|
		46	+\| 5 \| `ascendc_add5` \| `stream()` \| 等待 queue 完成后启动 \| 等价方式2 \|
36		47
37	-#### 封装Python模块	48	+> 推荐方式3：`stream(false)` 配合 `OpCommand::RunOpApiV2`，保留 TaskQueue 流水线性能，与 TorchNPU 内置算子行为一致。
38		49
39	-在`pybind11.asc`文件中使用了pybind11库来将C++代码封装成Python模块，在Python侧可以通过`import`方式进行调用。例如：	50	+### 2. Python 模块绑定
40	-
41	- ```c++
42	- PYBIND11_MODULE(custom_ops, m)
43	- {
44	- m.def("custom_add", &ascendc_ops::run_ascendc_add, "");
45	- }
46	- ```
47		51
48	- 通过此绑定，python侧可通过`op_extension.ops.custom_add`调用自定义的API。	52	+在 `pybind11.asc` 中使用 pybind11 将 C++ 函数暴露为 Python 接口。Python 侧仅暴露推荐方式3：
49		53
50	-#### Aten IR实现	54	+```cpp
		55	+namespace ascendc_ops {
		56	+// 方式3(推荐): stream(false) + OpCommand::RunOpApiV2
		57	+at::Tensor ascendc_add3(const at::Tensor &x, const at::Tensor &y);
		58	+at::Tensor run_trig_custom(const at::Tensor &x, const at::Tensor &out_sin, const at::Tensor &out_cos);
		59	+}
51		60
52	-根据Aten IR定义适配算子。	61	+PYBIND11_MODULE(custom_ops_lib, m)
53	-TorchNPU的算子下发和执行是异步的，通过TASKQUEUE实现，	62	+{
54	-样例中，我们通过`at_npu::native::OpCommand::RunOpApiV2`方法，将算子执行入队到TorchNPU的TASKQUEUE。样例如下：	63	+ m.def("custom_add", &ascendc_ops::ascendc_add3, "");
		64	+ m.def("custom_trig", &ascendc_ops::run_trig_custom, "");
		65	+}
		66	+```
55		67
56	- ```c++	68	+### 3. Aten IR 实现
		69	+
		70	+算子通过 `at_npu::native::OpCommand::RunOpApiV2` 入队到 TorchNPU 的 TaskQueue，实现异步下发。推荐方式3的实现：
		71	+
		72	+```cpp
57	#include "torch_npu/csrc/core/npu/NPUStream.h"	73	#include "torch_npu/csrc/core/npu/NPUStream.h"
58	#include "torch_npu/csrc/framework/OpCommand.h"	74	#include "torch_npu/csrc/framework/OpCommand.h"
		75	+
59	namespace ascendc_ops {	76	namespace ascendc_ops {
60	-at::Tensor run_ascendc_add(const at::Tensor &x, const at::Tensor &y)	77	+at::Tensor ascendc_add3(const at::Tensor &x, const at::Tensor &y)
61	{	78	{
62	- auto acl_stream = c10_npu::getCurrentNPUStream().stream(true);	79	+ // stream(false) 返回 ACL stream 但不清 queue
		80	+ auto acl_stream = c10_npu::getCurrentNPUStream().stream(false);
63	at::Tensor z = at::empty_like(x);	81	at::Tensor z = at::empty_like(x);
64	uint32_t blockDim = 8;	82	uint32_t blockDim = 8;
65	uint32_t totalLength = 1;	83	uint32_t totalLength = 1;
	@@ -67,20 +85,21 @@ at::Tensor run_ascendc_add(const at::Tensor &x, const at::Tensor &y)
67	totalLength *= size;	85	totalLength *= size;
68	}	86	}
69	// Launch the custom kernel use <<<>>>	87	// Launch the custom kernel use <<<>>>
70	- auto acl_call = [=]() -> int{	88	+ auto acl_call = [=]() -> int {
71	- add_custom<<<blockDim, nullptr, acl_stream>>>((uint8_t )(x.mutable_data_ptr()), (uint8_t )(y.mutable_data_ptr()),	89	+ add_custom<<<blockDim, nullptr, acl_stream>>>(
72	- (uint8_t *)(z.mutable_data_ptr()), totalLength);	90	+ (uint8_t *)(x.mutable_data_ptr()),
		91	+ (uint8_t *)(y.mutable_data_ptr()),
		92	+ (uint8_t *)(z.mutable_data_ptr()),
		93	+ totalLength);
73	return 0;	94	return 0;
74	};	95	};
75	at_npu::native::OpCommand::RunOpApiV2("ascendc_add", acl_call);	96	at_npu::native::OpCommand::RunOpApiV2("ascendc_add", acl_call);
76	-
77	return z;	97	return z;
78	}	98	}
79	-
80	} // namespace ascendc_ops	99	} // namespace ascendc_ops
81	- ```	100	+```
82		101
83	-上述主要介绍了自定义算子kernel集成的必备流程。	102	+上述主要介绍了自定义算子kernel集成的必备流程。完整的5种正确启动方式见`./csrc/add_custom.asc`，Python侧通过`./csrc/pybind11.asc`仅暴露推荐方式3。
84		103
85	最后，通过创建ops路径，定义python接口，通过`module_name.ops.custom_add`可以调用自定义算子。测试样例如下：	104	最后，通过创建ops路径，定义python接口，通过`module_name.ops.custom_add`可以调用自定义算子。测试样例如下：
86		105
	@@ -96,18 +115,31 @@ output = op_extension.ops.custom_add(x_npu, y_npu)
96		115
97	## 运行自定义的算子	116	## 运行自定义的算子
98		117
99	- 运行依赖PyTorch、TorchNPU和CANN。具体安装步骤参考[TorchNPU文档](https://gitcode.com/ascend/pytorch#%E5%AE%89%E8%A3%85)	118	+### 1. 编译 whl 包
100	- 运行流程：
101	-
102	- 1. 运行setup脚本，编译生成whl包。
103		119
104	```bash	120	```bash
105	python setup.py bdist_wheel	121	python setup.py bdist_wheel
106	```	122	```
107		123
108	- 我们的编译工程通过setuptools已为用户封装好如何编译算子kernel和集成到PyTorch。	124	+`setup.py` 关键逻辑：
109		125
110	- 2. 安装whl包	126	+\| 步骤 \| 实现 \| 说明 \|
		127	+\|------\|------\|------\|
		128	+\| 源码收集 \| `glob.glob("csrc/*.asc")` \| 自动收集所有 `.asc` 文件 \|
		129	+\| 架构识别 \| `get_npu_arch()` \| 通过 `npu-smi info` 解析芯片型号，映射到 `dav-2201`/`dav-3510` \|
		130	+\| 依赖路径 \| `get_dependency_paths()` \| 自动收集 torch / torch_npu / Python 的 include 与 lib 路径 \|
		131	+\| ABI 对齐 \| `torch._C._GLIBCXX_USE_CXX11_ABI` \| 与 PyTorch ABI 保持一致 \|
		132	+\| 编译器 \| `bisheng -x asc` \| 使用 CANN 提供的 bisheng 编译器 \|
		133	+
		134	+编译产物位于 `dist/op_extension-0.1-*.whl`。
		135	+
		136	+可选环境变量：
		137	+
		138	+```bash
		139	+USE_NINJA=1 python setup.py bdist_wheel # 启用 ninja 加速
		140	+```
		141	+
		142	+### 2. 安装 whl 包
111		143
112	```bash	144	```bash
113	cd dist	145	cd dist

Mexamples/cpp_extension/csrc/add_custom.asc+123-23

	@@ -135,43 +135,143 @@ __global__ __vector__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, uint32_t
135	}	135	}
136		136
137	// ------------------------------ C++接口函数 ------------------------------	137	// ------------------------------ C++接口函数 ------------------------------
138	-// ascendc_add: PyTorch和自定义内核之间的接口函数	138	+// 本文件提供 5 种正确的内核启动方式：
139	-// 参数：	139	+// - 方式1: 直接传入 NPUStream 对象，<<<>>> 内部清 queue
140	-// x, y: 输入张量	140	+// - 方式2: stream(true) 获取 ACL stream，返回前清 queue
141	-// 返回值：	141	+// - 方式3(推荐): stream(false) 配合 OpCommand，通过入 queue 方式同步
142	-// 加法结果张量	142	+// - 方式4: stream(true) 清 queue + OpCommand 入 queue
		143	+// - 方式5: stream() 获取 ACL stream，返回前等待 queue 内操作完成（等价方式2）
		144	+//
		145	+// 推荐使用方式3：性能最优（保留 TaskQueue 流水线），且与 TorchNPU 内置算子行为一致。
		146	+// Python 侧仅暴露方式3。
		147	+
143	namespace ascendc_ops {	148	namespace ascendc_ops {
144	-at::Tensor run_ascendc_add(const at::Tensor &x, const at::Tensor &y)	149	+
		150	+// 方式1: 清 queue - 直接传入 NPUStream 对象
		151	+// 使用 getCurrentNPUStream() 获取 NPUStream 对象，<<<>>> 内部会在内核启动前清 queue，
		152	+// 确保与之前任务的同步。
		153	+at::Tensor ascendc_add1(const at::Tensor &x, const at::Tensor &y)
145	{	154	{
146	- // 获取当前NPU流	155	+ // 获取 NPUStream 对象，此时不会清 queue
147	- auto acl_stream = c10_npu::getCurrentNPUStream().stream(true);	156	+ auto npu_stream = c10_npu::getCurrentNPUStream();
148	-
149	- // 创建与输入形状相同的空张量作为输出
150	at::Tensor z = at::empty_like(x);	157	at::Tensor z = at::empty_like(x);
151	-
152	- // 设置块维度
153	uint32_t blockDim = 8;	158	uint32_t blockDim = 8;
154	-
155	- // 计算输入张量的总元素数量
156	uint32_t totalLength = 1;	159	uint32_t totalLength = 1;
157	for (uint32_t size : x.sizes()) {	160	for (uint32_t size : x.sizes()) {
158	totalLength *= size;	161	totalLength *= size;
159	}	162	}
		163	+ auto xGm = (uint8_t *)(x.mutable_data_ptr());
		164	+ auto yGm = (uint8_t *)(y.mutable_data_ptr());
		165	+ auto zGm = (uint8_t *)(z.mutable_data_ptr());
		166	+ // <<<>>> 传入 NPUStream 对象，内部会在内核启动前清 queue
		167	+ add_custom<<<blockDim, nullptr, npu_stream>>>(xGm, yGm, zGm, totalLength);
		168	+ return z;
		169	+}
160		170
161	- // 定义内核启动lambda函数	171	+
162	- auto acl_call = [=]() -> int{	172	+// 方式2: 清 queue - 使用 stream(true) 获取 ACL stream 并清 queue
163	- // 启动自定义内核	173	+// stream(true) 返回 ACL stream(aclrtStream)，在返回前会清 queue，确保与之前任务的同步。
164	- add_custom<<<blockDim, nullptr, acl_stream>>>((uint8_t *)(x.mutable_data_ptr()),	174	+// 与方式1等价。
165	- (uint8_t *)(y.mutable_data_ptr()),	175	+at::Tensor ascendc_add2(const at::Tensor &x, const at::Tensor &y)
166	- (uint8_t *)(z.mutable_data_ptr()),	176	+{
167	- totalLength);	177	+ // stream(true) 在返回 ACL stream 前会清 queue
		178	+ auto acl_stream = c10_npu::getCurrentNPUStream().stream(true);
		179	+ at::Tensor z = at::empty_like(x);
		180	+ uint32_t blockDim = 8;
		181	+ uint32_t totalLength = 1;
		182	+ for (uint32_t size : x.sizes()) {
		183	+ totalLength *= size;
		184	+ }
		185	+ auto xGm = (uint8_t *)(x.mutable_data_ptr());
		186	+ auto yGm = (uint8_t *)(y.mutable_data_ptr());
		187	+ auto zGm = (uint8_t *)(z.mutable_data_ptr());
		188	+ // Launch kernel use <<<>>>
		189	+ add_custom<<<blockDim, nullptr, acl_stream>>>(xGm, yGm, zGm, totalLength);
		190	+ return z;
		191	+}
		192	+
		193	+
		194	+// 方式3: Lambda 入 queue - 使用 stream(false) 配合 OpCommand 进行 queue 管理 [推荐]
		195	+// stream(false) 返回 ACL stream 但不清 queue，结合 OpCommand::RunOpApiV2 使用，
		196	+// 内核启动被封装在 lambda 中，通过正确地入 queue 和出 queue 确保正确的执行顺序，
		197	+// 同时保留 TaskQueue 的流水线性能。
		198	+at::Tensor ascendc_add3(const at::Tensor &x, const at::Tensor &y)
		199	+{
		200	+ // stream(false) 返回 ACL stream 但不清 queue
		201	+ auto acl_stream = c10_npu::getCurrentNPUStream().stream(false);
		202	+ at::Tensor z = at::empty_like(x);
		203	+ uint32_t blockDim = 8;
		204	+ uint32_t totalLength = 1;
		205	+ for (uint32_t size : x.sizes()) {
		206	+ totalLength *= size;
		207	+ }
		208	+ auto xGm = (uint8_t *)(x.mutable_data_ptr());
		209	+ auto yGm = (uint8_t *)(y.mutable_data_ptr());
		210	+ auto zGm = (uint8_t *)(z.mutable_data_ptr());
		211	+
		212	+ // 定义内核启动 lambda 函数
		213	+ auto acl_call = [=]() -> int {
		214	+ // Launch kernel use <<<>>>
		215	+ add_custom<<<blockDim, nullptr, acl_stream>>>(xGm, yGm, zGm, totalLength);
168	return 0;	216	return 0;
169	};	217	};
170		218
171	- // 通过OpCommand运行内核	219	+ // 通过 OpCommand 运行内核，内部会进行入 queue 和出 queue 的操作，
		220	+ // 确保与 stream 中其他 NPU 操作的正确同步
172	at_npu::native::OpCommand::RunOpApiV2("ascendc_add", acl_call);	221	at_npu::native::OpCommand::RunOpApiV2("ascendc_add", acl_call);
		222	+ return z;
		223	+}
173		224
174	- // 返回计算结果	225	+
		226	+// 方式4: 清 queue + Lambda 入 queue
		227	+// stream(true) 返回 ACL stream(aclrtStream) 并清 queue，然后在 lambda 中使用该 stream。
		228	+// 这种组合是安全的：先清 queue 确保之前的任务完成，再通过 OpCommand 入 queue 执行内核。
		229	+// 与方式3相比会损失 TaskQueue 流水线性能，但语义更直观。
		230	+at::Tensor ascendc_add4(const at::Tensor &x, const at::Tensor &y)
		231	+{
		232	+ // stream(true) 在返回 ACL stream 前会清 queue
		233	+ auto acl_stream = c10_npu::getCurrentNPUStream().stream(true);
		234	+ at::Tensor z = at::empty_like(x);
		235	+ uint32_t blockDim = 8;
		236	+ uint32_t totalLength = 1;
		237	+ for (uint32_t size : x.sizes()) {
		238	+ totalLength *= size;
		239	+ }
		240	+ auto xGm = (uint8_t *)(x.mutable_data_ptr());
		241	+ auto yGm = (uint8_t *)(y.mutable_data_ptr());
		242	+ auto zGm = (uint8_t *)(z.mutable_data_ptr());
		243	+
		244	+ // 定义内核启动 lambda 函数
		245	+ auto acl_call = [=]() -> int {
		246	+ // Launch kernel use <<<>>>
		247	+ add_custom<<<blockDim, nullptr, acl_stream>>>(xGm, yGm, zGm, totalLength);
		248	+ return 0;
		249	+ };
		250	+
		251	+ // 通过 OpCommand 运行内核，内部会进行入 queue 和出 queue 的操作
		252	+ at_npu::native::OpCommand::RunOpApiV2("ascendc_add", acl_call);
		253	+ return z;
		254	+}
		255	+
		256	+
		257	+// 方式5: 清 queue - 使用 stream() 获取 ACL stream 并等待 queue 内操作
		258	+// stream()（无参数）在返回 ACL stream(aclrtStream) 前会等待 queue 内操作完成，
		259	+// 与方式2(stream(true))等价，确保与之前任务的同步。
		260	+at::Tensor ascendc_add5(const at::Tensor &x, const at::Tensor &y)
		261	+{
		262	+ // stream() 无参数，返回 ACL stream 前会等待 queue 内操作完成
		263	+ auto acl_stream = c10_npu::getCurrentNPUStream().stream();
		264	+ at::Tensor z = at::empty_like(x);
		265	+ uint32_t blockDim = 8;
		266	+ uint32_t totalLength = 1;
		267	+ for (uint32_t size : x.sizes()) {
		268	+ totalLength *= size;
		269	+ }
		270	+ auto xGm = (uint8_t *)(x.mutable_data_ptr());
		271	+ auto yGm = (uint8_t *)(y.mutable_data_ptr());
		272	+ auto zGm = (uint8_t *)(z.mutable_data_ptr());
		273	+ // Launch kernel use <<<>>>
		274	+ add_custom<<<blockDim, nullptr, acl_stream>>>(xGm, yGm, zGm, totalLength);
175	return z;	275	return z;
176	}	276	}
177		277

Mexamples/cpp_extension/csrc/pybind11.asc+4-2

	@@ -2,13 +2,15 @@
2	#include <torch/extension.h>	2	#include <torch/extension.h>
3		3
4	namespace ascendc_ops {	4	namespace ascendc_ops {
5	-at::Tensor run_ascendc_add(const at::Tensor &x, const at::Tensor &y);	5	+// 方式3(推荐): stream(false) + OpCommand::RunOpApiV2，保留 TaskQueue 流水线性能
		6	+at::Tensor ascendc_add3(const at::Tensor &x, const at::Tensor &y);
6	at::Tensor run_trig_custom(const at::Tensor &x, const at::Tensor &out_sin, const at::Tensor &out_cos);	7	at::Tensor run_trig_custom(const at::Tensor &x, const at::Tensor &out_sin, const at::Tensor &out_cos);
7	}	8	}
8		9
9	// expose Ascend custom ops to Python	10	// expose Ascend custom ops to Python
10	PYBIND11_MODULE(custom_ops_lib, m)	11	PYBIND11_MODULE(custom_ops_lib, m)
11	{	12	{
12	- m.def("custom_add", &ascendc_ops::run_ascendc_add, "");	13	+ // Python 侧仅暴露推荐方式3，其他方式仅作为源码参考
		14	+ m.def("custom_add", &ascendc_ops::ascendc_add3, "");
13	m.def("custom_trig", &ascendc_ops::run_trig_custom, "");	15	m.def("custom_trig", &ascendc_ops::run_trig_custom, "");
14	}	16	}

Mexamples/cpp_extension/csrc/trig_inplace_custom.asc+15-4

	@@ -14,6 +14,7 @@
14	#include <torch/extension.h>	14	#include <torch/extension.h>
15		15
16	#include "torch_npu/csrc/core/npu/NPUStream.h"	16	#include "torch_npu/csrc/core/npu/NPUStream.h"
		17	+#include "torch_npu/csrc/framework/OpCommand.h"
17	#include "kernel_operator.h"	18	#include "kernel_operator.h"
18		19
19	constexpr uint32_t BUFFER_NUM = 2; //tensor num for each queue	20	constexpr uint32_t BUFFER_NUM = 2; //tensor num for each queue
	@@ -106,9 +107,11 @@ __global__ __vector__ void trig_inplace_custom(GM_ADDR x, GM_ADDR out_sin, GM_AD
106	}	107	}
107		108
108	namespace ascendc_ops {	109	namespace ascendc_ops {
		110	+// 使用推荐方式3: stream(false) + OpCommand::RunOpApiV2，保留 TaskQueue 流水线性能
109	at::Tensor run_trig_custom(const at::Tensor &x, const at::Tensor &out_sin, const at::Tensor &out_cos)	111	at::Tensor run_trig_custom(const at::Tensor &x, const at::Tensor &out_sin, const at::Tensor &out_cos)
110	{	112	{
111	- auto acl_stream = c10_npu::getCurrentNPUStream().stream(true);	113	+ // stream(false) 返回 ACL stream 但不清 queue
		114	+ auto acl_stream = c10_npu::getCurrentNPUStream().stream(false);
112	at::Tensor out_tan = at::empty_like(x);	115	at::Tensor out_tan = at::empty_like(x);
113	uint32_t blockDim = 8;	116	uint32_t blockDim = 8;
114	uint32_t totalLength = 1;	117	uint32_t totalLength = 1;
	@@ -119,8 +122,16 @@ at::Tensor run_trig_custom(const at::Tensor &x, const at::Tensor &out_sin, const
119	auto sinGm = static_cast<uint8_t >(const_cast<void >(out_sin.storage().data()));	122	auto sinGm = static_cast<uint8_t >(const_cast<void >(out_sin.storage().data()));
120	auto cosGm = static_cast<uint8_t >(const_cast<void >(out_cos.storage().data()));	123	auto cosGm = static_cast<uint8_t >(const_cast<void >(out_cos.storage().data()));
121	auto tanGm = static_cast<uint8_t >(const_cast<void >(out_tan.storage().data()));	124	auto tanGm = static_cast<uint8_t >(const_cast<void >(out_tan.storage().data()));
122	- // Launch the custom kernel using <<<>>>	125	+
123	- trig_inplace_custom<<<blockDim, nullptr, acl_stream>>>(xGm, sinGm, cosGm, tanGm, totalLength);	126	+ // 定义内核启动 lambda 函数
		127	+ auto acl_call = [=]() -> int {
		128	+ // Launch the custom kernel using <<<>>>
		129	+ trig_inplace_custom<<<blockDim, nullptr, acl_stream>>>(xGm, sinGm, cosGm, tanGm, totalLength);
		130	+ return 0;
		131	+ };
		132	+
		133	+ // 通过 OpCommand 运行内核，确保与 stream 中其他 NPU 操作的正确同步
		134	+ at_npu::native::OpCommand::RunOpApiV2("ascendc_trig", acl_call);
124	return out_tan;	135	return out_tan;
125	}	136	}
126	-} // namespace ascendc_ops	137	+} // namespace ascendc_ops