添加aclgraph+<<<>>> 示例代码-op-plugin-AtomGit

已合并

添加aclgraph+<<<>>> 示例代码 #4019

mihudan创建于 1月9日从mihudan/op-plugin： master 合入到 Ascend/op-plugin： master

添加aclgraph+<<<>>> 示例代码 #4019

已合并

mihudan创建于 1月9日

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

mihudan

1月9日

What type of PR is this?

Uncomment only one /kind <> line, hit enter to put that in a new line, and remove leading whitespaces from that line:

/kind bug
/kind task
/kind feature

What does this PR do / why do we need it:
添加aclgraph 与<<<>>>联调的demo测试代码，进行验证并供用户参考。

展示了如何使用PyTorch的torch.librar以及pybind两张方式注册自定义算子，通过<<<>>>内核调用符调用核函数，并适配aclgraph使用该自定义算子，以简单的Add算子和三角函数计算的原地算子为例，实现aclgraph下自定义算子的调用。

展示了3种aclgraph的使能方式，通过对比NPU输出与CPU标准加法结果来验证自定义算子的数值正确性。

torch.npu.NPUGraph()
torch.npu.make_graphed_callables
backend="npugraph_ex"

算子描述

Add算子

算子功能：
Add算子实现了两个数据相加，返回相加结果的功能。对应的算子原型为：
```
ascendc_add(Tensor x, Tensor y) -> Tensor
```
算子规格：

核函数名 add_custom

算子输入 name shape data type format

x 8 * 2048 int ND

y 8 * 2048 int ND

算子输出 z 8 * 2048 int ND

原地三角函数算子

算子功能：
该算子入参为x, out_sin ,out_cos, 算子调用后，out_sin会被原地修改为sin(x)计算结果，out_cos会被原地修改为cos(x)计算结果，返回值tan(x)计算结果。对应的算子原型为：
```
ascendc_trig(Tensor x, Tensor(a!) out_sin, Tensor(b!) out_cos) -> Tensor
```

算子规格：

核函数名	trig_inplace_custom
算子输入	name	shape	data type	format
	x	8 * 2048	float	ND
	out_sin	8 * 2048	float	ND
	out_cos	8 * 2048	float	ND
算子输出	out_sin	8 * 2048	float	ND
	out_cos	8 * 2048	float	ND
	out_tan	8 * 2048	float	ND

代码实现介绍

以Add算子为例，样例在*.asc文件中定义了一个名为ascendc_ops的命名空间，并在其中注册了ascendc_add函数。在ascendc_add函数中通过c10_npu::getCurrentNPUStream()函数获取当前NPU上的流，并通过内核调用符<<<>>>调用自定义的Kernel函数add_custom，在NPU上执行算子。
```
  add_custom<<<blockDim, nullptr, aclStream>>>(xGm, yGm, zGm, totalLength);
```
PyTorch提供TORCH_LIBRARY_FRAGMENT宏作为自定义算子注册的核心接口，用于创建并初始化自定义算子库，注册后在Python侧可以通过torch.ops.namespace.op_name方式进行调用，例如：
```
TORCH_LIBRARY_FRAGMENT(ascendc_ops, m)
{
    m.def(ascendc_add"(Tensor x, Tensor y) -> Tensor");
}
```
TORCH_LIBRARY_IMPL用于将算子逻辑绑定到特定的DispatchKey（PyTorch设备调度标识）。针对NPU设备，需要将算子实现注册到PrivateUse1这一专属的DispatchKey上，例如：
```
TORCH_LIBRARY_IMPL(ascendc_ops, PrivateUse1, m)
{
    m.impl("ascendc_add", TORCH_FN(ascendc_ops::ascendc_add));
}
```
注册Meta函数：

注册Meta函数使faketensor流程正常工作，在使用fx, compile等功能涉及，注册代码如下：
```
TORCH_LIBRARY_IMPL(ascendc_ops, Meta, m)
{
  m.impl("ascendc_add", &add_impl_meta);
}
```

Special notes for your reviewers:

Pull Request已成功合入, 合并人@ascend-robot

（感谢 mihudan 的贡献）

M mihudan

1月9日创建了 pull request，commit 4b024d18

ascend-robot成员

1月9日评论：

Thank your for your pull-request.

The full list of commands accepted by me can be found at here.

You can get sig-info at here

AtlasAccount成员

1月9日评论：

openLiBing可接管PR, PR详情链接:
https://www.openlibing.com/apps/transitionPrDetail?owner=Ascend&repo=op-plugin&number=4019&platform=gitcode&codeHostingPlatformFlag=gitcode

ascend-robot成员

1月9日评论：

以下是根据您提交的修改文件推荐的Reviewer和Committer序列，需各模块评审通过后方可合入

Module List	Reviewers	Committers
repo-Ascend/op-plugin	qingfenxiaochong1, dilililiwhy, wangchao430, wang-guangbin, yangkaixin	liwei386, dilililiwhy, wangchao430, yanpengquan, yangkaixin

ascend-robot成员

1月9日添加了label：ascend-cla/yes

此处折叠了182条消息查看更多

yang-binrong成员1月17日进行代码检视1

examples/kernel_extension_aclgraph/torch_library/test/trig_aclgraph_test.py

54+ # Test using make_graphed_callables

55+ def test_npugraph_ex_backend(self):

56+ model = Model().npu()

57+ compiled_model = torch.compile(model, backend="npugraph_ex", fullgraph=True, dynamic=True)

yang-binrong1月17日评论：

fullgrap需要与out_tan风格一致

yang-binrong成员1月17日进行代码检视1

examples/kernel_extension_aclgraph/torch_library/test/trig_aclgraph_test.py

11+ return out_tan

12+

13+

14+length = [8, 2048]

yang-binrong1月17日评论：

需要以dtype相关sizeof计算替换2048,防止功能拓展时散弹式修改

yang-binrong成员1月17日进行代码检视1

examples/kernel_extension_aclgraph/torch_library/csrc/trig_inplace_custom.asc

109+{

110+ auto acl_stream = c10_npu::getCurrentNPUStream().stream(true);

111+ at::Tensor out_tan = at::empty_like(x);

112+ uint32_t blockDim = 8;

yang-binrong1月17日评论：

需要以dtype相关sizeof计算替换8与1,防止功能拓展时散弹式修改

yang-binrong成员1月17日进行代码检视1

examples/kernel_extension_aclgraph/pybind/test/add_aclgraph_test.py

85+ x, y = self.get_rand_input()

86+ output = torch.ops.ascendc_ops.ascendc_add(x.npu(), y.npu()).cpu()

87+ cpuout = torch.add(x, y)

88+ self.assertEqual(output, cpuout)

yang-binrong1月17日评论：

cpuout需要改为cpu_out

yang-binrong成员1月17日进行代码检视1

examples/kernel_extension_aclgraph/pybind/csrc/add_custom.asc

93+{

94+ auto acl_stream = c10_npu::getCurrentNPUStream().stream(true);

95+ at::Tensor z = at::empty_like(x);

96+ uint32_t blockDim = 8;

yang-binrong1月17日评论：

需要以dtype相关sizeof计算替换数字,防止功能拓展时散弹式修改

核函数名	add_custom
算子输入	name	shape	data type	format
	x	8 * 2048	int	ND
	y	8 * 2048	int	ND
算子输出	z	8 * 2048	int	ND