Star241
45
代码
代码
Issues68
Pull Requests51
Actions
项目讨论
Wiki
分析
Star241
45
已合并
chore: update repo metadata and compliance config #56
LiuZonggu创建于 6月29日
chore: update repo metadata and compliance config #56
已合并
LiuZonggu创建于 6月29日
从已删除 :repo-metadata-compliance合入到cann/ops-rasmaster
共 16 个文件变更+226-747
M.clang-format+9-8
@@ -1,9 +1,9 @@
1-# This program is free software, you can redistribute it and/or modify.1+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.2+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3-# This file is a part of the CANN Open Software.3+# CANN Open Software License Agreement Version 2.0 (the "License").
4-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
5# Please refer to the License for details. You may not use this file except in compliance with the License.4# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.5+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7# See LICENSE in the root of the software repository for the full text of the License.7# See LICENSE in the root of the software repository for the full text of the License.
8# ============================================================================8# ============================================================================
9 9 
@@ -22,7 +22,8 @@ BraceWrapping:
22 AfterCaseLabel: false22 AfterCaseLabel: false
23 AfterClass: false23 AfterClass: false
24 AfterControlStatement: false24 AfterControlStatement: false
25- AfterEnum: true25+ # AfterEnum: false (紧凑模式，左大括号和enum同行, enum Color {...), AfterEnum: true (左大括号和enum换行).
26+ AfterEnum: false
26 AfterFunction: true27 AfterFunction: true
27 AfterNamespace: false28 AfterNamespace: false
28 AfterStruct: false29 AfterStruct: false
@@ -42,7 +43,7 @@ AlignOperands: true
42AlignTrailingComments: true43AlignTrailingComments: true
43AllowAllArgumentsOnNextLine: true44AllowAllArgumentsOnNextLine: true
44AllowAllParametersOfDeclarationOnNextLine: true45AllowAllParametersOfDeclarationOnNextLine: true
45-AlwaysBreakBeforeMultilineStrings: true46+AlwaysBreakBeforeMultilineStrings: false
46AlwaysBreakTemplateDeclarations: true47AlwaysBreakTemplateDeclarations: true
47 48 
48BinPackArguments: true49BinPackArguments: true
@@ -56,7 +57,7 @@ PointerAlignment: Left
56# single line layout57# single line layout
57AllowShortBlocksOnASingleLine: false58AllowShortBlocksOnASingleLine: false
58AllowShortCaseLabelsOnASingleLine: false59AllowShortCaseLabelsOnASingleLine: false
59-AllowShortFunctionsOnASingleLine: false60+AllowShortFunctionsOnASingleLine: true
60AllowShortIfStatementsOnASingleLine: false61AllowShortIfStatementsOnASingleLine: false
61AllowShortLoopsOnASingleLine: false62AllowShortLoopsOnASingleLine: false
62 63 
A.gitignore+76-0
@@ -0,0 +1,76 @@
1+# binary files
2+*.so
3+ 
4+cmake-build-debugindocker/
5+.DS_Store
6+# distribution and packaging
7+build/
8+build_out/
9+output/
10+dist/
11+*.egg-info/
12+*.egg
13+*.whl
14+/third_party
15+ 
16+# vim cache
17+*.swp
18+ 
19+# python cache
20+*.pyc
21+__pycache__
22+MANIFEST.in
23+ 
24+# python ut/st
25+cov_report/
26+report/
27+data/
28+*.dump
29+!**/st/msopgen/simulator/*.dump
30+tf_plugin_model.pb
31+V*_onnx_plugin_model.onnx
32+ 
33+# protobuf
34+*.pb.h
35+*.pb.c
36+ 
37+# cann tmp
38+lib/
39+kernel_meta/
40+ 
41+# Jetbrain
42+.idea
43+cmake-build-debug
44+cmake-build-release
45+ 
46+# Visual Studio
47+.vs
48+ 
49+# Visual Studio Code
50+.vscode
51+*.tmp
52+*.temp
53+ 
54+# ModelArts
55+.arts
56+ 
57+# UT report file
58+ut_report.txt
59+ 
60+# ascend package download dir
61+ascend_download
62+ 
63+# makeself
64+makeself
65+ 
66+# tbetoolkits log files
67+tbetoolkits-*.log
68+*_result.csv
69+ 
70+# local build package
71+single/
72+scripts/ci/tmp/
73+pr_filelist.txt
74+oat_reports/
75+ 
76+log/
A.pre-commit-config.yaml+76-0
@@ -0,0 +1,76 @@
1+# ---------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ---------------------------------------------------------------------------------------------------------
10+minimum_pre_commit_version: 4.0.0
11+exclude: ^LICENSES/|\.(html|csv|svg)$
12+ 
13+default_stages: [pre-commit]
14+ci:
15+ autofix_prs: false
16+ autoupdate_schedule: monthly
17+ 
18+repos:
19+ # pre-commit 基础检查
20+ - repo: https://gitcode.com/pre-commit/pre-commit-hooks
21+ rev: v4.6.0
22+ hooks:
23+ - id: trailing-whitespace
24+ - id: end-of-file-fixer
25+ - id: check-yaml
26+ args: ["--allow-multiple-documents"]
27+ - id: check-added-large-files
28+ - id: check-merge-conflict
29+ - id: detect-private-key
30+ - id: check-json
31+ 
32+ # C++ 代码格式化检查
33+ - repo: https://gitcode.com/pre-commit-clang/mirrors-clang-format
34+ rev: v18.1.8
35+ hooks:
36+ - id: clang-format
37+ files: \.(c|h|cpp|hpp|cc|hh|cxx|hxx|asc)$
38+ args:
39+ - "--style=file"
40+ - "--verbose"
41+ - "-i"
42+ exclude: ^build/|tests/third_party/
43+ 
44+ # Python 代码检查
45+ - repo: https://gitcode.com/gh_mirrors/ru/ruff-pre-commit
46+ rev: v0.14.14
47+ hooks:
48+ - id: ruff-check
49+ args: ["--output-format", "github", "--fix"]
50+ types: [python]
51+ - id: ruff-format
52+ types: [python]
53+ 
54+ # 拼写检查
55+ - repo: https://gitcode.com/gh_mirrors/co/codespell
56+ rev: v2.4.1
57+ hooks:
58+ - id: codespell
59+ args: [
60+ "-L",
61+ "CANN,cann,NNAL,nnal,ASCEND,ascend,EnQue,CopyIn,ArchType,AND,ND,tbe,copyin,alog",
62+ "--skip",
63+ "*.py,*.cpp,*.hpp,*.c,*.h",
64+ ]
65+ 
66+ # OAT 合规性检查 (Open Source Audit Tool)
67+ # - repo: local
68+ # hooks:
69+ # - id: oat-check
70+ # name: OAT Compliance Check
71+ # entry: bash scripts/oat_check.sh
72+ # language: system
73+ # pass_filenames: true
74+ # types: [file]
75+ # stages: [pre-commit]
76+ # verbose: true
ACHANGELOG.md+1-0
@@ -0,0 +1 @@
1+# Changelog
MOAT.xml+10-8
@@ -1,10 +1,10 @@
1<?xml version="1.0" encoding="UTF-8"?>1<?xml version="1.0" encoding="UTF-8"?>
2-<!-- Copyright (c) 2025 Huawei Technologies Co., Ltd.2+<!-- Copyright (c) 2026 Huawei Technologies Co., Ltd.
3- This program is free software, you can redistribute it and/or modify it under the terms and conditions of 3+ This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4 CANN Open Software License Agreement Version 2.0 (the "License").4 CANN Open Software License Agreement Version 2.0 (the "License").
5 Please refer to the License for details. You may not use this file except in compliance with the License.5 Please refer to the License for details. You may not use this file except in compliance with the License.
6- THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+ THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. 7+ INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8 See LICENSE in the root of the software repository for the full text of the License.8 See LICENSE in the root of the software repository for the full text of the License.
9-->9-->
10 10 
@@ -14,7 +14,7 @@
14 <policylist>14 <policylist>
15 <policy name="projectPolicy" desc="">15 <policy name="projectPolicy" desc="">
16 <!--policyitem type="compatibility" name="GPL-2.0+" path="abc/.*" desc="Process that runs independently, invoked by the X process."/-->16 <!--policyitem type="compatibility" name="GPL-2.0+" path="abc/.*" desc="Process that runs independently, invoked by the X process."/-->
17- <policyitem type="license" name="cann License" path=".*" desc=""/>17+ <policyitem type="license" name="CANN-2.0" path=".*" desc=""/>
18 <policyitem type="copyright" name="Huawei Technologies Co., Ltd." path=".*" rule="may" group="defaultGroup" filefilter="copyrightPolicyFilter" desc=""/>18 <policyitem type="copyright" name="Huawei Technologies Co., Ltd." path=".*" rule="may" group="defaultGroup" filefilter="copyrightPolicyFilter" desc=""/>
19 </policy>19 </policy>
20 </policylist>20 </policylist>
@@ -28,6 +28,9 @@
28 <!--filteritem type="filename" name="*.uvwxyz" desc="Describe the reason for filtering scan results"/-->28 <!--filteritem type="filename" name="*.uvwxyz" desc="Describe the reason for filtering scan results"/-->
29 <!--filteritem type="filepath" name="abcdefg/.*.uvwxyz" desc="Describe the reason for filtering scan results"/-->29 <!--filteritem type="filepath" name="abcdefg/.*.uvwxyz" desc="Describe the reason for filtering scan results"/-->
30 <!--filteritem type="filepath" name="projectroot/[a-zA-Z0-9]{20,}.sh" desc="Temp files"/-->30 <!--filteritem type="filepath" name="projectroot/[a-zA-Z0-9]{20,}.sh" desc="Temp files"/-->
31+ <!-- Exclude YAML config files -->
32+ <filteritem type="filename" name="*.yaml" desc="YAML config files"/>
33+ <filteritem type="filename" name="*.yml" desc="YAML config files"/>
31 </filefilter>34 </filefilter>
32 <filefilter name="copyrightPolicyFilter" desc="Filters for copyright header policies">35 <filefilter name="copyrightPolicyFilter" desc="Filters for copyright header policies">
33 <!--filteritem type="filename" name="*.uvwxyz" desc="Describe the reason for filtering scan results"/-->36 <!--filteritem type="filename" name="*.uvwxyz" desc="Describe the reason for filtering scan results"/-->
@@ -59,12 +62,11 @@
59 <licensematcherlist>62 <licensematcherlist>
60 <licensematcher name="cann License" desc="If the scanning result is InvalidLicense, you can define matching rules here. Note that quotation marks must be escaped.">63 <licensematcher name="cann License" desc="If the scanning result is InvalidLicense, you can define matching rules here. Note that quotation marks must be escaped.">
61 <licensetext name="64 <licensetext name="
62- Copyright (c) 2025 Huawei Technologies Co., Ltd.
63 This program is free software, you can redistribute it and/or modify it under the terms and conditions of 65 This program is free software, you can redistribute it and/or modify it under the terms and conditions of
64 CANN Open Software License Agreement Version 2.0 (the &quot;License&quot;).66 CANN Open Software License Agreement Version 2.0 (the &quot;License&quot;).
65 Please refer to the License for details. You may not use this file except in compliance with the License.67 Please refer to the License for details. You may not use this file except in compliance with the License.
66- THIS SOFTWARE IS PROVIDED ON AN &quot;AS IS&quot; BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 68+ THIS SOFTWARE IS PROVIDED ON AN &quot;AS IS&quot; BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
67- INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. 69+ INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
68 See LICENSE in the root of the software repository for the full text of the License.70 See LICENSE in the root of the software repository for the full text of the License.
69 " desc=""/>71 " desc=""/>
70 </licensematcher>72 </licensematcher>
MSECURITY.md+13-11
@@ -7,9 +7,9 @@
7## 文件权限控制7## 文件权限控制
8 8 
9- 建议用户在主机（包括宿主机）及容器中设置运行系统umask值为0027及以上，保障新增文件夹默认最高权限为750，新增文件默认最高权限为640。9- 建议用户在主机（包括宿主机）及容器中设置运行系统umask值为0027及以上，保障新增文件夹默认最高权限为750，新增文件默认最高权限为640。
10-- 建议用户对个人隐私数据、商业资产、源文件和算子开发过程中保存的各类文件等敏感内容做好权限控制等安全措施。例如涉及本项目安装目录权限管控、输入公共数据文件权限管控，设定的权限建议参考[A-文件（夹）各场景权限管控推荐最大值](#A-文件（夹）各场景权限管控推荐最大值)。10+- 建议用户对个人隐私数据、商业资产、源文件和算子开发过程中保存的各类文件等敏感内容做好权限控制等安全措施。例如涉及本项目安装目录权限管控、输入公共数据文件权限管控，设定的权限建议参考[A-文件（夹）各场景权限管控推荐最大值](#a-文件夹各场景权限管控推荐最大值)。
11- 算子运行时可能会缓存算子编译文件，存储在运行目录下的`kernel_meta_*`文件夹内，加快后续算子的调用速度，用户可根据需要自行对生成后的相关文件进行权限控制。11- 算子运行时可能会缓存算子编译文件，存储在运行目录下的`kernel_meta_*`文件夹内，加快后续算子的调用速度，用户可根据需要自行对生成后的相关文件进行权限控制。
12-- 用户安装和使用过程需要做好权限控制，建议参考[A-文件（夹）各场景权限管控推荐最大值](#A-文件（夹）各场景权限管控推荐最大值)文件权限参考进行设置。12+- 用户安装和使用过程需要做好权限控制，建议参考[A-文件（夹）各场景权限管控推荐最大值](#a-文件夹各场景权限管控推荐最大值)文件权限参考进行设置。
13 13 
14## 构建安全声明14## 构建安全声明
15 15 
@@ -22,20 +22,22 @@
22- 算子通过[PyTorch](https://gitee.com/ascend/pytorch)方式调用时，可能会因为版本不匹配导致运行错误，具体请参考[PyTorch安全声明](https://gitee.com/ascend/pytorch#%E5%AE%89%E5%85%A8%E5%A3%B0%E6%98%8E)。22- 算子通过[PyTorch](https://gitee.com/ascend/pytorch)方式调用时，可能会因为版本不匹配导致运行错误，具体请参考[PyTorch安全声明](https://gitee.com/ascend/pytorch#%E5%AE%89%E5%85%A8%E5%A3%B0%E6%98%8E)。
23 23 
24## 公网地址声明24## 公网地址声明
25+ 
25本项目代码中包含的公网地址声明如下所示：26本项目代码中包含的公网地址声明如下所示：
26 27 
27| 类型 | 开源代码地址 | 文件名 | 公网IP地址/公网URL地址/域名/邮箱地址/压缩文件地址 | 用途说明 |28| 类型 | 开源代码地址 | 文件名 | 公网IP地址/公网URL地址/域名/邮箱地址/压缩文件地址 | 用途说明 |
28| :------------: |:------------------------------------------------------------------------------------------:|:----------------------------------------------------------| :---------------------------------------------------------- |:-----------------------------------------|29| :------------: |:------------------------------------------------------------------------------------------:|:----------------------------------------------------------| :---------------------------------------------------------- |:-----------------------------------------|
29-| 依赖 | 不涉及 | cmake/third_party/makeself-fetch.cmake | https://gitcode.com/cann-src-third-party/makeself/releases/download/release-2.5.0-patch1.0/makeself-release-2.5.0-patch1.tar.gz | 从gitcode下载makeself源码，作用编译依赖 |30+| 依赖 | 不涉及 | cmake/third_party/makeself-fetch.cmake | [https://gitcode.com/cann-src-third-party/makeself/releases/download/release-2.5.0-patch1.0/makeself-release-2.5.0-patch1.tar.gz](https://gitcode.com/cann-src-third-party/makeself/releases/download/release-2.5.0-patch1.0/makeself-release-2.5.0-patch1.tar.gz) | 从gitcode下载makeself源码，作用编译依赖 |
30-| 依赖 | 不涉及 | cmake/third_party/nlohmann_json.cmake | https://gitcode.com/cann-src-third-party/json/releases/download/v3.11.3/include.zip | 从gitcode下载json源码，作用编译依赖 |31+| 依赖 | 不涉及 | cmake/third_party/nlohmann_json.cmake | [https://gitcode.com/cann-src-third-party/json/releases/download/v3.11.3/include.zip](https://gitcode.com/cann-src-third-party/json/releases/download/v3.11.3/include.zip) | 从gitcode下载json源码，作用编译依赖 |
31-| 依赖 | 不涉及 | cmake/third_party/gtest.cmake | https://gitcode.com/cann-src-third-party/googletest/releases/download/v1.14.0/googletest-1.14.0.tar.gz | 从gitcode下载googletest源码，作用编译依赖 |32+| 依赖 | 不涉及 | cmake/third_party/gtest.cmake | [https://gitcode.com/cann-src-third-party/googletest/releases/download/v1.14.0/googletest-1.14.0.tar.gz](https://gitcode.com/cann-src-third-party/googletest/releases/download/v1.14.0/googletest-1.14.0.tar.gz) | 从gitcode下载googletest源码，作用编译依赖 |
32-| 依赖 | 不涉及 | cmake/third_party/eigen.cmake | https://gitcode.com/cann-src-third-party/eigen/releases/download/3.4.0/eigen-3.4.0.tar.gz | 从gitcode下载eigen源码，作用编译依赖 |33+| 依赖 | 不涉及 | cmake/third_party/eigen.cmake | [https://gitcode.com/cann-src-third-party/eigen/releases/download/5.0.0-h0.trunk/eigen-5.0.0.tar.gz](https://gitcode.com/cann-src-third-party/eigen/releases/download/5.0.0-h0.trunk/eigen-5.0.0.tar.gz) | 从gitcode下载eigen源码，作用编译依赖 |
33-| 依赖 | 不涉及 | ops-ras/install_deps.sh | https://apt.kitware.com/keys/kitware-archive-latest.asc|从gitcode下载install_deps源码，作用编译依赖 |34+| 依赖 | 不涉及 | ops-ras/install_deps.sh | [https://apt.kitware.com/keys/kitware-archive-latest.asc](https://apt.kitware.com/keys/kitware-archive-latest.asc) | 从gitcode下载install_deps源码，作用编译依赖 |
34-| 依赖 | 不涉及 | ops-ras/install_deps.sh | https://apt.kitware.com/ubuntu/|从gitcode下载install_deps源码，作用编译依赖 |35+| 依赖 | 不涉及 | ops-ras/install_deps.sh | [https://apt.kitware.com/ubuntu/](https://apt.kitware.com/ubuntu/) | 从gitcode下载install_deps源码，作用编译依赖 |
35-| 依赖 | 不涉及 | cmake | https://apt.kitware.com/keys/kitware-archive-latest.asc | 从kitware下载cmake软件，作用编译依赖 |36+| 依赖 | 不涉及 | cmake | [https://apt.kitware.com/keys/kitware-archive-latest.asc](https://apt.kitware.com/keys/kitware-archive-latest.asc) | 从kitware下载cmake软件，作用编译依赖 |
36-| 依赖 | 不涉及 | cmake | https://apt.kitware.com/ubuntu/ | 从kitware下载cmake软件，作用编译依赖 |37+| 依赖 | 不涉及 | cmake | [https://apt.kitware.com/ubuntu/](https://apt.kitware.com/ubuntu/) | 从kitware下载cmake软件，作用编译依赖 |
37 38 
38## 漏洞机制说明39## 漏洞机制说明
40+ 
39[漏洞管理](https://gitcode.com/cann/community/blob/master/security/security.md)41[漏洞管理](https://gitcode.com/cann/community/blob/master/security/security.md)
40 42 
41## 附录43## 附录
@@ -60,4 +62,4 @@
60| 业务数据文件目录 | 750（rwxr-x---） |62| 业务数据文件目录 | 750（rwxr-x---） |
61| 密钥组件、私钥、证书、密文文件目录 | 700（rwx-----） |63| 密钥组件、私钥、证书、密文文件目录 | 700（rwx-----） |
62| 密钥组件、私钥、证书、加密密文 | 600（rw-------） |64| 密钥组件、私钥、证书、加密密文 | 600（rw-------） |
63-| 加解密接口、加解密脚本 | 500（r-x------） |65+| 加解密接口、加解密脚本 | 500（r-x------） |
AThird_Party_Open_Source_Software_List.yaml+20-0
@@ -0,0 +1,20 @@
1+#(全量组件，type缺省为run，存在测试组件类型test，构建组件类型build，辅助类插件plugin.以及发行组件类型run，需要额外类型联系安全sig)
2+#YAML格式，靠空格缩进表达层级关系（禁止用Tab缩进），缩进2个空格（也可以是4个，只要同级元素缩进一致即可）
3+abseil-cpp:
4+ version: 20230802.1
5+ type: plugin
6+googletest:
7+ version: v1.14.0
8+ type: test
9+eigen:
10+ version: 3.4.0
11+ type: run
12+makeself:
13+ version: 2.5.0-patch1.0
14+ type: build
15+json:
16+ version: v3.11.3
17+ type: run
18+protobuf:
19+ version: v25.1
20+ type: run
MThird_Party_Open_Source_Software_Notice+16-16
@@ -1,9 +1,9 @@
1-* This program is free software, you can redistribute it and/or modify.1+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
2-* Copyright (c) 2025 Huawei Technologies Co., Ltd.2+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
3-* This file is a part of the CANN Open Software.3+* CANN Open Software License Agreement Version 2.0 (the "License").
4-* Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
5* Please refer to the License for details. You may not use this file except in compliance with the License.4* Please refer to the License for details. You may not use this file except in compliance with the License.
6-* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.5+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
6+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
7* See LICENSE in the root of the software repository for the full text of the License.7* See LICENSE in the root of the software repository for the full text of the License.
8 8 
9OPEN SOURCE SOFTWARE NOTICE9OPEN SOURCE SOFTWARE NOTICE
@@ -62,8 +62,8 @@ Software: makeself release-2.5.0-patch1
62Copyright notice:62Copyright notice:
63Copyright (c) 1989, 1991 Free Software Foundation, Inc. copyrighted by the Free Software Foundation Copyright (c) 1998-2023 by Stephane Peter megastep@megastep.org63Copyright (c) 1989, 1991 Free Software Foundation, Inc. copyrighted by the Free Software Foundation Copyright (c) 1998-2023 by Stephane Peter megastep@megastep.org
64 64 
65-License: GPL V2.0 65+License: GPL V2.0
66-GNU GENERAL PUBLIC LICENSE 66+GNU GENERAL PUBLIC LICENSE
67Version 2, June 199167Version 2, June 1991
68 68 
69Copyright (C) 1989, 1991 Free Software Foundation, Inc.69Copyright (C) 1989, 1991 Free Software Foundation, Inc.
@@ -85,7 +85,7 @@ Finally, any free program is threatened constantly by software patents. We wish
85 85 
86The precise terms and conditions for copying, distribution and modification follow.86The precise terms and conditions for copying, distribution and modification follow.
87 87 
88-TERMS AND CONDITIONS FOR COPYING, DISTRIBUTION AND MODIFICATION 88+TERMS AND CONDITIONS FOR COPYING, DISTRIBUTION AND MODIFICATION
890. This License applies to any program or other work which contains a notice placed by the copyright holder saying it may be distributed under the terms of this General Public License. The "Program", below, refers to any such program or work, and a "work based on the Program" means either the Program or any derivative work under copyright law: that is to say, a work containing the Program or a portion of it, either verbatim or with modifications and/or translated into another language. (Hereinafter, translation is included without limitation in the term "modification".) Each licensee is addressed as "you".890. This License applies to any program or other work which contains a notice placed by the copyright holder saying it may be distributed under the terms of this General Public License. The "Program", below, refers to any such program or work, and a "work based on the Program" means either the Program or any derivative work under copyright law: that is to say, a work containing the Program or a portion of it, either verbatim or with modifications and/or translated into another language. (Hereinafter, translation is included without limitation in the term "modification".) Each licensee is addressed as "you".
90 90 
91Activities other than copying, distribution and modification are not covered by this License; they are outside its scope. The act of running the Program is not restricted, and the output from the Program is covered only if its contents constitute a work based on the Program (independent of having been made by running the Program). Whether that is true depends on what the Program does.91Activities other than copying, distribution and modification are not covered by this License; they are outside its scope. The act of running the Program is not restricted, and the output from the Program is covered only if its contents constitute a work based on the Program (independent of having been made by running the Program). Whether that is true depends on what the Program does.
@@ -222,7 +222,7 @@ THE SOFTWARE.
222 222 
223Copyright Notice and License Texts223Copyright Notice and License Texts
224Software: eigen 3.4.0224Software: eigen 3.4.0
225-Copyright notice: 225+Copyright notice:
226Copyright (c) 2020 Arm Ltd. and Contributors226Copyright (c) 2020 Arm Ltd. and Contributors
227Copyright (c) 2011, Intel Corporation227Copyright (c) 2011, Intel Corporation
228Copyright (c) 2007 Free Software Foundation, Inc. http://fsf.org/228Copyright (c) 2007 Free Software Foundation, Inc. http://fsf.org/
@@ -1466,7 +1466,7 @@ Mozilla Public License Version 2.0
1466 means any form of the work other than Source Code Form.1466 means any form of the work other than Source Code Form.
1467 1467 
14681.7. "Larger Work"14681.7. "Larger Work"
1469- means a work that combines Covered Software with other material, in 1469+ means a work that combines Covered Software with other material, in
1470 a separate file or files, that is not Covered Software.1470 a separate file or files, that is not Covered Software.
1471 1471 
14721.8. "License"14721.8. "License"
@@ -1814,7 +1814,7 @@ Usage: "
1814/control/assert/op_graph/assert_proto.h,1814/control/assert/op_graph/assert_proto.h,
1815/pooling/avg_pool3_d/op_graph/avg_pool3_d_proto.h"1815/pooling/avg_pool3_d/op_graph/avg_pool3_d_proto.h"
1816 1816 
1817-Copyright Notice: 1817+Copyright Notice:
1818Copyright 2019-2020 Huawei Technologies Co., Ltd1818Copyright 2019-2020 Huawei Technologies Co., Ltd
1819 1819 
1820License: Apache License V2.01820License: Apache License V2.0
@@ -1829,7 +1829,7 @@ Usage: "
1829/tests/ut/op_kernel/data_utils.cpp1829/tests/ut/op_kernel/data_utils.cpp
1830/scripts/kernel/binary_script/parser_ini.py"1830/scripts/kernel/binary_script/parser_ini.py"
1831 1831 
1832-Copyright Notice: 1832+Copyright Notice:
1833Copyright 2019-2020 Huawei Technologies Co., Ltd1833Copyright 2019-2020 Huawei Technologies Co., Ltd
1834 1834 
1835License: Apache License V2.01835License: Apache License V2.0
@@ -1843,7 +1843,7 @@ Usage: "
1843/examples/add_example_aicpu/op_kernel_aicpu/add_example_aicpu.cpp1843/examples/add_example_aicpu/op_kernel_aicpu/add_example_aicpu.cpp
1844/examples/add_example_aicpu/op_kernel_aicpu/add_example_aicpu.h"1844/examples/add_example_aicpu/op_kernel_aicpu/add_example_aicpu.h"
1845 1845 
1846-Copyright Notice: 1846+Copyright Notice:
1847Copyright 2019-2020 Huawei Technologies Co., Ltd1847Copyright 2019-2020 Huawei Technologies Co., Ltd
1848 1848 
1849License: Apache License V2.01849License: Apache License V2.0
@@ -1856,7 +1856,7 @@ Official Repository: https://github.com/Ascend/samples
1856Usage: "1856Usage: "
1857/scripts/custom/install.sh"1857/scripts/custom/install.sh"
1858 1858 
1859-Copyright Notice: 1859+Copyright Notice:
1860Copyright 2019-2020 Huawei Technologies Co., Ltd1860Copyright 2019-2020 Huawei Technologies Co., Ltd
1861 1861 
1862License: Apache License V2.01862License: Apache License V2.0
@@ -1871,7 +1871,7 @@ Usage: "
1871/conv/conv3d_transpose_v2/op_graph/conv3d_transpose_v2_proto.h1871/conv/conv3d_transpose_v2/op_graph/conv3d_transpose_v2_proto.h
1872/rnn/dynamic_rnn/op_graph/dynamic_rnn_proto.h"1872/rnn/dynamic_rnn/op_graph/dynamic_rnn_proto.h"
1873 1873 
1874-Copyright Notice: 1874+Copyright Notice:
1875Copyright 2019-2020 Huawei Technologies Co., Ltd1875Copyright 2019-2020 Huawei Technologies Co., Ltd
1876 1876 
1877License: Apache License V2.01877License: Apache License V2.0
@@ -1885,7 +1885,7 @@ Usage: "
1885/scripts/util/insert_op_info.py1885/scripts/util/insert_op_info.py
1886/scripts/util/parse_ini_to_json.py"1886/scripts/util/parse_ini_to_json.py"
1887 1887 
1888-Copyright Notice: 1888+Copyright Notice:
1889Copyright 2019-2020 Huawei Technologies Co., Ltd1889Copyright 2019-2020 Huawei Technologies Co., Ltd
1890 1890 
1891License: Apache License V2.01891License: Apache License V2.0
Mclassify_rule.yaml+4-8
@@ -1,5 +1,5 @@
1# ---------------------------------------------------------------------------------------------------------1# ---------------------------------------------------------------------------------------------------------
2-# Copyright (c) 2025-2026 Huawei Technologies Co., Ltd.2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3# This program is free software, you can redistribute it and/or modify it under the terms and conditions of3# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
5# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -43,15 +43,11 @@ ras@ops-ras:
43 kernel_style: null43 kernel_style: null
44 unrelease:44 unrelease:
45 test_code:45 test_code:
46- - ops/ops-ras/example/46+ - ops/ops-ras/examples/
47- - ops/ops-ras/example/add_example_aicpu/examples47+ - ops/ops-ras/experimental/
48- - ops/ops-ras/example/add_example_aicpu/tests
49- - ops/ops-ras/example/add_example/examples
50- - ops/ops-ras/example/add_example/tests
51 - ops/ops-ras/tests/48 - ops/ops-ras/tests/
52 - ops/ops-ras/build.sh49 - ops/ops-ras/build.sh
53 - ops/ops-ras/CMakeLists.txt50 - ops/ops-ras/CMakeLists.txt
54- - ops/ops-ras/common/inc/op_api/aclnn_check.h
55 - ops/ops-ras/common/inc/op_api/level2_base.h51 - ops/ops-ras/common/inc/op_api/level2_base.h
56 non_delivery: null52 non_delivery: null
57 open_source: null53 open_source: null
@@ -78,4 +74,4 @@ reliability@ops-ras:
78 - ops/ops-ras/reliability/obfuscation_setup/tests/74 - ops/ops-ras/reliability/obfuscation_setup/tests/
79 - ops/ops-ras/reliability/crypto/tests/75 - ops/ops-ras/reliability/crypto/tests/
80 non_delivery: null76 non_delivery: null
81- open_source: null77+ open_source: null
Dcmake/Third_Party_Open_Source_Software_List.yaml+0-21
@@ -1,21 +0,0 @@
1-observability:
2- googletest:
3- cpeName: googletest
4- groupId:
5- url: https://github.com/google/googletest.git
6- version: v1.14.0
7- eigen:
8- cpeName: eigen
9- groupId:
10- url: https://github.com/libeigen/eigen.git
11- version: 3.4.0
12- makeself:
13- cpeName: makeself
14- groupId:
15- url: https://github.com/megastep/makeself.git
16- version: release-2.5.0-patch1
17- json:
18- cpeName: json
19- groupId:
20- url: https://github.com/nlohmann/json.git
21- version: v3.11.3
Mexamples/fast_kernel_launch_example/ascend_ops/__init__.py+1-4
@@ -1,7 +1,7 @@
1#!/usr/bin/env python31#!/usr/bin/env python3
2# -*- coding: utf-8 -*-2# -*- coding: utf-8 -*-
3# ----------------------------------------------------------------------------3# ----------------------------------------------------------------------------
4-# Copyright (c) 2025 Huawei Technologies Co., Ltd.4+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
5# This program is free software, you can redistribute it and/or modify it under the terms and conditions of5# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
6# CANN Open Software License Agreement Version 2.0 (the "License").6# CANN Open Software License Agreement Version 2.0 (the "License").
7# Please refer to the License for details. You may not use this file except in compliance with the License.7# Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -12,7 +12,6 @@
12 12 
13"""Ascend Ops"""13"""Ascend Ops"""
14__version__ = "1.0.0"14__version__ = "1.0.0"
15-import torch
16 15 
17try:16try:
18 from . import _C17 from . import _C
@@ -20,5 +19,3 @@ except ImportError as e:
20 raise ImportError(19 raise ImportError(
21 "Cannot import _C. Please make sure the `ascend_ops` is properly installed. "20 "Cannot import _C. Please make sure the `ascend_ops` is properly installed. "
22 ) from e21 ) from e
23- 
24-from . import ops
Dexamples/fast_kernel_launch_example/ascend_ops/ops.py+0-34
@@ -1,34 +0,0 @@
1-#!/usr/bin/env python3
2-# -*- coding: UTF-8 -*-
3-# ----------------------------------------------------------------------------
4-# This program is free software, you can redistribute it and/or modify it.
5-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6-# This file is a part of the CANN Open Software.
7-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
8-# Please refer to the License for details. You may not use this file except in compliance with the License.
9-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
10-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
11-# See LICENSE in the root of the software repository for the full text of the License.
12-# ----------------------------------------------------------------------------
13- 
14-import torch
15-from torch import Tensor
16-from ascend_ops.trans_conv3d_format import ncdhw_to_ndc1hwc0, ncdhw_to_fz3d, ndc1hwc0_to_ncdhw
17- 
18-__all__ = ["conv3d_custom", ]
19- 
20- 
21-def conv3d_custom(input: Tensor, weight: Tensor, strides: list, pads: list, dilations: list,
22- bias: Tensor = None, enable_hf32: bool = False) -> Tensor:
23- print(torch.ops.ascend_ops.conv3d_custom)
24- assert hasattr(torch.ops.ascend_ops, "conv3d_custom"), "The 'conv3d_custom' operator is not registered in the 'torch.ops.ascend_ops' namespace."
25-
26- origin_input_shape = list(input.shape)
27- origin_weight_shape = list(weight.shape)
28- origin_cout = origin_weight_shape[0]
29- input = ncdhw_to_ndc1hwc0(input).npu()
30- weight = ncdhw_to_fz3d(weight).npu()
31- bias = bias.npu() if bias is not None else None
32- output = torch.ops.ascend_ops.conv3d_custom(input, weight, strides, pads, dilations, origin_input_shape,
33- origin_weight_shape, bias, enable_hf32).cpu()
34- return ndc1hwc0_to_ncdhw(output, origin_cout)
Dexamples/fast_kernel_launch_example/ascend_ops/trans_conv3d_format.py+0-91
@@ -1,91 +0,0 @@
1-#!/usr/bin/env python3
2-# -*- coding: UTF-8 -*-
3-# ----------------------------------------------------------------------------
4-# This program is free software, you can redistribute it and/or modify it.
5-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6-# This file is a part of the CANN Open Software.
7-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
8-# Please refer to the License for details. You may not use this file except in compliance with the License.
9-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
10-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
11-# See LICENSE in the root of the software repository for the full text of the License.
12-# ----------------------------------------------------------------------------
13- 
14-import math
15-import torch
16- 
17-CUBE_N = 16
18- 
19-K0_MAP = {
20- torch.float32: 8,
21- torch.float16: 16,
22- torch.bfloat16: 16,
23-}
24- 
25- 
26-def lcm(m, n):
27- return (m * n) // math.gcd(m, n)
28- 
29- 
30-def ceil(m, n):
31- return (m + n - 1) // n
32- 
33- 
34-def ncdhw_to_ndc1hwc0(input_tensor):
35- cube_k = K0_MAP[input_tensor.dtype]
36- input_shape = input_tensor.shape
37- dim_c = input_shape[1]
38- pad_value = (cube_k - dim_c % cube_k) % cube_k
39- res_tensor = torch.nn.functional.pad(input_tensor, (0, 0, 0, 0, 0, 0, 0, pad_value), mode="constant", value=0)
40- dim_c1 = (dim_c + pad_value) // cube_k
41- res_tensor = res_tensor.reshape((input_shape[0], dim_c1, cube_k, input_shape[2], input_shape[3], input_shape[4]))
42- res_tensor = res_tensor.permute(0, 3, 1, 4, 5, 2)
43- return res_tensor.contiguous()
44- 
45- 
46-def ndc1hwc0_to_ncdhw(output_tensor, real_cout):
47- def ndc1hwc0_to_ndchw():
48- output_shape = output_tensor.shape
49- res_shape = [output_shape[0], output_shape[1], real_cout, output_shape[3], output_shape[4]]
50- res_tensor = torch.zeros(res_shape, dtype=output_tensor.dtype)
51- for i in range(output_shape[0]):
52- for j in range(output_shape[1]):
53- for k in range(output_shape[2]):
54- for m in range(output_shape[5]):
55- if k * output_shape[5] + m < real_cout:
56- res_tensor[i, j, k * output_shape[5] + m, :, :] = output_tensor[i, j, k, :, :, m]
57- return res_tensor
58- 
59- res_tensor = ndc1hwc0_to_ndchw()
60- res_tensor = res_tensor.permute(0, 2, 1, 3, 4)
61- return res_tensor.contiguous()
62- 
63- 
64-def ncdhw_to_fz3d(filter_tensor, groups=1):
65- cube_k = K0_MAP[filter_tensor.dtype]
66- 
67- cout, cin_ori, dk, kh, kw = filter_tensor.shape
68- cout_ori = cout // groups
69- enlarge = min(lcm(lcm(cin_ori, cube_k) // cin_ori, lcm(cout_ori, CUBE_N) // cout_ori), groups)
70- cin_opt = ceil(enlarge * cin_ori, cube_k) * cube_k
71- cout_opt = ceil(enlarge * cout_ori, CUBE_N) * CUBE_N
72- group_opt = ceil(groups, enlarge)
73- 
74- filter_shape_gdc1hwnc0 = [group_opt, dk, cin_opt // cube_k, kh, kw, cout_opt, cube_k]
75- res_tensor = torch.zeros(filter_shape_gdc1hwnc0, dtype=filter_tensor.dtype)
76- for g in range(groups):
77- for ci in range(cin_ori):
78- for co in range(cout_ori):
79- for kernel_d in range(dk):
80- e = g % enlarge
81- dst_ci = e * cin_ori + ci
82- dst_co = e * cout_ori + co
83- src_co = g * cout_ori + co
84- res_tensor[g // enlarge, kernel_d, dst_ci // cube_k, :, :, dst_co,
85- dst_ci % cube_k] = filter_tensor[
86- src_co, ci,
87- kernel_d, :,
88- :]
89- res_tensor = res_tensor.reshape(group_opt * dk * (cin_opt // cube_k) * kh * kw, cout_opt // CUBE_N, CUBE_N,
90- cube_k)
91- return res_tensor.contiguous()
Dexamples/fast_kernel_launch_example/csrc/conv3d_custom/ascend910b/CMakeLists.txt+0-24
@@ -1,24 +0,0 @@
1-# ----------------------------------------------------------------------------
2-# This program is free software, you can redistribute it and/or modify it.
3-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
4-# This file is a part of the CANN Open Software.
5-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6-# Please refer to the License for details. You may not use this file except in compliance with the License.
7-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9-# See LICENSE in the root of the software repository for the full text of the License.
10-# ----------------------------------------------------------------------------
11- 
12-set(CONV3D_TILING_DIR "${CMAKE_CURRENT_SOURCE_DIR}/../../../../../conv/conv3d_v2/op_host/op_tiling")
13-set(CONV3D_TILING_SOURCES
14- "${CONV3D_TILING_DIR}/conv3d_api_tiling_algorithm.cpp"
15- "${CONV3D_TILING_DIR}/conv3d_api_tiling_algorithm_hw_mode.cpp"
16- "${CONV3D_TILING_DIR}/conv3d_api_tiling_algorithm_pointwise.cpp"
17- "${CONV3D_TILING_DIR}/conv3d_api_tiling_base.cpp"
18- "${CONV3D_TILING_DIR}/conv3d_api_tiling.cpp"
19- "${CONV3D_TILING_DIR}/conv3d_api_tiling_utils.cpp"
20- "${CONV3D_TILING_DIR}/conv3d_tiling_engine.cpp"
21- "${CONV3D_TILING_DIR}/conv3d_tiling_utils.cpp"
22-)
23- 
24-add_sources("--npu-arch=dav-2201" "${CONV3D_TILING_SOURCES}")
Dexamples/fast_kernel_launch_example/csrc/conv3d_custom/ascend910b/conv3d_custom_torch.cpp+0-429
@@ -1,429 +0,0 @@
1-/**
2- * This program is free software, you can redistribute it and/or modify it.
3- * Copyright (c) 2025 Huawei Technologies Co., Ltd.
4- * This file is a part of the CANN Open Software.
5- * Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
6- * Please refer to the License for details. You may not use this file except in compliance with the License.
7- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
8- * BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
9- * See LICENSE in the root of the software repository for the full text of the License.
10- */
11- 
12-/*!
13- * \file conv3d_custom_torch.cpp
14- * \brief
15- */
16- 
17-#include <ATen/Operators.h>
18-#include <torch/all.h>
19-#include <torch/library.h>
20-#include "acl/acl.h"
21-#include "torch_npu/csrc/core/npu/NPUStream.h"
22-#include "torch_npu/csrc/core/npu/DeviceUtils.h"
23-#include "torch_npu/csrc/framework/OpCommand.h"
24-#include "tiling/platform/platform_ascendc.h"
25-#include "kernel_operator.h"
26- 
27-#include "conv/conv3d_v2/op_host/op_tiling/conv3d_tiling_engine.h"
28-#include "conv/conv3d_v2/op_kernel/conv3d_v2_kernel_dispatch.h"
29- 
30-namespace ascend_ops {
31- 
32-namespace Conv3dCustom {
33-using namespace std;
34-using namespace Ops::Ras::Conv3dV2;
35-using namespace Ops::Ras::Conv3dTilingEngineApi;
36- 
37-constexpr static uint8_t NCDHW_DIM = 5;
38-constexpr static uint8_t NDC1HWC0_DIM = 6;
39-constexpr static uint8_t FRACTAL_Z_3D_DIM = 4;
40-constexpr static uint8_t CONV_ATTRS_DIM = 3;
41-constexpr static uint8_t N_DIM_IDX_NCDHW = 0;
42-constexpr static uint8_t C_DIM_IDX_NCDHW = 1;
43-constexpr static uint8_t D_DIM_IDX_NCDHW = 2;
44-constexpr static uint8_t H_DIM_IDX_NCDHW = 3;
45-constexpr static uint8_t W_DIM_IDX_NCDHW = 4;
46-constexpr static uint8_t N_DIM_IDX_NDC1HWC0 = 0;
47-constexpr static uint8_t D_DIM_IDX_NDC1HWC0 = 1;
48-constexpr static uint8_t H_DIM_IDX_NDC1HWC0 = 3;
49-constexpr static uint8_t W_DIM_IDX_NDC1HWC0 = 4;
50-constexpr static uint8_t ATTRS_D_DIM_IDX_NCDHW = 0;
51-constexpr static uint8_t ATTRS_H_DIM_IDX_NCDHW = 1;
52-constexpr static uint8_t ATTRS_W_DIM_IDX_NCDHW = 2;
53- 
54-constexpr static uint8_t PBUFFERFLAG_L0A_DB_MASK = 1;
55-constexpr static uint8_t PBUFFERFLAG_L0B_DB_MASK = 2;
56- 
57-static std::map<at::ScalarType, uint8_t> K0_BY_DTYPE_MAP = {{at::kHalf, 16}, {at::kBFloat16, 16}, {at::kFloat, 8}};
58- 
59-template <typename T>
60-struct ToAscendDtype {
61- using Type = float;
62-};
63- 
64-template <>
65-struct ToAscendDtype<c10::BFloat16> {
66- using Type = bfloat16_t;
67-};
68- 
69-template <>
70-struct ToAscendDtype<c10::Half> {
71- using Type = half;
72-};
73- 
74-template <typename T>
75-struct ToAscendBiasDtype {
76- using Type = float;
77-};
78- 
79-template <>
80-struct ToAscendBiasDtype<c10::Half> {
81- using Type = half;
82-};
83- 
84-template <typename InputType, typename WeightType, typename OutputType, typename BiasType>
85-__aicore__ void L0PingPongAllClose(
86- __gm__ uint8_t* input, __gm__ uint8_t* weight, __gm__ uint8_t* bias, __gm__ uint8_t* output,
87- const Conv3DV2TilingData& tilingData)
88-{
89- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_OFF) &&
90- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::M_Mode)) {
91- DispatchConv3DV2Kernel<
92- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::ALL_CLOSE, ConvBL1ByPass::BYPASS_OFF,
93- GroupConvType::NoGroup_Conv, OutputOrder::M_Mode, QuantType::NO_QUANT>(
94- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
95- return;
96- }
97- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_OFF) &&
98- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::HW_Mode)) {
99- DispatchConv3DV2Kernel<
100- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::ALL_CLOSE, ConvBL1ByPass::BYPASS_OFF,
101- GroupConvType::NoGroup_Conv, OutputOrder::HW_Mode, QuantType::NO_QUANT>(
102- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
103- return;
104- }
105- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_ON) &&
106- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::M_Mode)) {
107- DispatchConv3DV2Kernel<
108- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::ALL_CLOSE, ConvBL1ByPass::BYPASS_ON,
109- GroupConvType::NoGroup_Conv, OutputOrder::M_Mode, QuantType::NO_QUANT>(
110- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
111- return;
112- }
113- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_ON) &&
114- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::HW_Mode)) {
115- DispatchConv3DV2Kernel<
116- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::ALL_CLOSE, ConvBL1ByPass::BYPASS_ON,
117- GroupConvType::NoGroup_Conv, OutputOrder::HW_Mode, QuantType::NO_QUANT>(
118- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
119- return;
120- }
121-}
122- 
123-template <typename InputType, typename WeightType, typename OutputType, typename BiasType>
124-__aicore__ void L0BPingPongOpen(
125- __gm__ uint8_t* input, __gm__ uint8_t* weight, __gm__ uint8_t* bias, __gm__ uint8_t* output,
126- const Conv3DV2TilingData& tilingData)
127-{
128- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_OFF) &&
129- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::M_Mode)) {
130- DispatchConv3DV2Kernel<
131- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::L0B_OPEN, ConvBL1ByPass::BYPASS_OFF,
132- GroupConvType::NoGroup_Conv, OutputOrder::M_Mode, QuantType::NO_QUANT>(
133- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
134- return;
135- }
136- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_OFF) &&
137- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::HW_Mode)) {
138- DispatchConv3DV2Kernel<
139- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::L0B_OPEN, ConvBL1ByPass::BYPASS_OFF,
140- GroupConvType::NoGroup_Conv, OutputOrder::HW_Mode, QuantType::NO_QUANT>(
141- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
142- return;
143- }
144- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_ON) &&
145- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::M_Mode)) {
146- DispatchConv3DV2Kernel<
147- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::L0B_OPEN, ConvBL1ByPass::BYPASS_ON,
148- GroupConvType::NoGroup_Conv, OutputOrder::M_Mode, QuantType::NO_QUANT>(
149- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
150- return;
151- }
152- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_ON) &&
153- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::HW_Mode)) {
154- DispatchConv3DV2Kernel<
155- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::L0B_OPEN, ConvBL1ByPass::BYPASS_ON,
156- GroupConvType::NoGroup_Conv, OutputOrder::HW_Mode, QuantType::NO_QUANT>(
157- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
158- return;
159- }
160-}
161- 
162-template <typename InputType, typename WeightType, typename OutputType, typename BiasType>
163-__aicore__ void L0APingPongOpen(
164- __gm__ uint8_t* input, __gm__ uint8_t* weight, __gm__ uint8_t* bias, __gm__ uint8_t* output,
165- const Conv3DV2TilingData& tilingData)
166-{
167- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_OFF) &&
168- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::M_Mode)) {
169- DispatchConv3DV2Kernel<
170- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::L0A_OPEN, ConvBL1ByPass::BYPASS_OFF,
171- GroupConvType::NoGroup_Conv, OutputOrder::M_Mode, QuantType::NO_QUANT>(
172- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
173- return;
174- }
175- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_OFF) &&
176- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::HW_Mode)) {
177- DispatchConv3DV2Kernel<
178- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::L0A_OPEN, ConvBL1ByPass::BYPASS_OFF,
179- GroupConvType::NoGroup_Conv, OutputOrder::HW_Mode, QuantType::NO_QUANT>(
180- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
181- return;
182- }
183- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_ON) &&
184- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::M_Mode)) {
185- DispatchConv3DV2Kernel<
186- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::L0A_OPEN, ConvBL1ByPass::BYPASS_ON,
187- GroupConvType::NoGroup_Conv, OutputOrder::M_Mode, QuantType::NO_QUANT>(
188- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
189- return;
190- }
191- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_ON) &&
192- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::HW_Mode)) {
193- DispatchConv3DV2Kernel<
194- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::L0A_OPEN, ConvBL1ByPass::BYPASS_ON,
195- GroupConvType::NoGroup_Conv, OutputOrder::HW_Mode, QuantType::NO_QUANT>(
196- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
197- return;
198- }
199-}
200- 
201-template <typename InputType, typename WeightType, typename OutputType, typename BiasType>
202-__aicore__ void L0PingPongAllOpen(
203- __gm__ uint8_t* input, __gm__ uint8_t* weight, __gm__ uint8_t* bias, __gm__ uint8_t* output,
204- const Conv3DV2TilingData& tilingData)
205-{
206- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_OFF) &&
207- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::M_Mode)) {
208- DispatchConv3DV2Kernel<
209- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::ALL_OPEN, ConvBL1ByPass::BYPASS_OFF,
210- GroupConvType::NoGroup_Conv, OutputOrder::M_Mode, QuantType::NO_QUANT>(
211- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
212- return;
213- }
214- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_OFF) &&
215- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::HW_Mode)) {
216- DispatchConv3DV2Kernel<
217- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::ALL_OPEN, ConvBL1ByPass::BYPASS_OFF,
218- GroupConvType::NoGroup_Conv, OutputOrder::HW_Mode, QuantType::NO_QUANT>(
219- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
220- return;
221- }
222- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_ON) &&
223- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::M_Mode)) {
224- DispatchConv3DV2Kernel<
225- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::ALL_OPEN, ConvBL1ByPass::BYPASS_ON,
226- GroupConvType::NoGroup_Conv, OutputOrder::M_Mode, QuantType::NO_QUANT>(
227- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
228- return;
229- }
230- if (tilingData.conv3dApiTiling.bl1BypassFlag == static_cast<uint8_t>(ConvBL1ByPass::BYPASS_ON) &&
231- tilingData.conv3dApiTiling.outputOrder == static_cast<int8_t>(OutputOrder::HW_Mode)) {
232- DispatchConv3DV2Kernel<
233- InputType, WeightType, OutputType, BiasType, ConvL0PingPong::ALL_OPEN, ConvBL1ByPass::BYPASS_ON,
234- GroupConvType::NoGroup_Conv, OutputOrder::HW_Mode, QuantType::NO_QUANT>(
235- input, weight, bias, nullptr, nullptr, output, nullptr, &tilingData);
236- return;
237- }
238-}
239- 
240-template <typename T>
241-__global__ __aicore__ void Conv3dCustomKernel(
242- __gm__ uint8_t* input, __gm__ uint8_t* weight, __gm__ uint8_t* bias, __gm__ uint8_t* output,
243- const Conv3DV2TilingData tilingData)
244-{
245- if constexpr (std::is_same_v<T, c10::Half> || std::is_same_v<T, c10::BFloat16> || std::is_same_v<T, float>) {
246- using InputType = ConvType<TPosition::GM, ConvFormat::NDC1HWC0, typename ToAscendDtype<T>::Type>;
247- using WeightType = ConvType<TPosition::GM, ConvFormat::FRACTAL_Z_3D, typename ToAscendDtype<T>::Type>;
248- using OutputType = ConvType<TPosition::GM, ConvFormat::NDC1HWC0, typename ToAscendDtype<T>::Type>;
249- using BiasType = ConvType<TPosition::GM, ConvFormat::ND, typename ToAscendBiasDtype<T>::Type>;
250- 
251- bool l0aDbFlag = tilingData.conv3dApiTiling.pBufferFlag & PBUFFERFLAG_L0A_DB_MASK;
252- bool l0bDbFlag = tilingData.conv3dApiTiling.pBufferFlag & PBUFFERFLAG_L0B_DB_MASK;
253- if (!l0aDbFlag && !l0bDbFlag) {
254- L0PingPongAllClose<InputType, WeightType, OutputType, BiasType>(input, weight, bias, output, tilingData);
255- return;
256- }
257- if (!l0aDbFlag && l0bDbFlag) {
258- L0BPingPongOpen<InputType, WeightType, OutputType, BiasType>(input, weight, bias, output, tilingData);
259- return;
260- }
261- if (l0aDbFlag && !l0bDbFlag) {
262- L0APingPongOpen<InputType, WeightType, OutputType, BiasType>(input, weight, bias, output, tilingData);
263- return;
264- }
265- if (l0aDbFlag && l0bDbFlag) {
266- L0PingPongAllOpen<InputType, WeightType, OutputType, BiasType>(input, weight, bias, output, tilingData);
267- return;
268- }
269- }
270-}
271- 
272-template <typename T>
273-void Conv3dCustomApi(
274- aclrtStream stream, const at::Tensor& input, const at::Tensor& weight, const torch::Tensor& output,
275- const vector<int64_t> strideList, const vector<int64_t> paddingList, const vector<int64_t> dilationList,
276- const vector<int64_t> oriInputShapeList, const vector<int64_t> oriWeightShapeList,
277- const c10::optional<torch::Tensor>& bias, bool enableHF32)
278-{
279- vector<int64_t> oriOutputShapeList = {
280- output.size(N_DIM_IDX_NDC1HWC0), oriWeightShapeList[N_DIM_IDX_NCDHW], output.size(D_DIM_IDX_NDC1HWC0),
281- output.size(H_DIM_IDX_NDC1HWC0), output.size(W_DIM_IDX_NDC1HWC0)};
282- 
283- __gm__ uint8_t* input_ptr = (__gm__ uint8_t*)input.data_ptr<T>();
284- __gm__ uint8_t* weight_ptr = (__gm__ uint8_t*)weight.data_ptr<T>();
285- __gm__ uint8_t* bias_ptr = nullptr;
286- __gm__ uint8_t* output_ptr = (__gm__ uint8_t*)output.data_ptr<T>();
287- Conv3dTilingEngine conv3dTilingEngine;
288- conv3dTilingEngine.SetOrgWeightShape(oriWeightShapeList);
289- conv3dTilingEngine.SetOrgFmapShape(oriInputShapeList);
290- conv3dTilingEngine.SetOrgOutputShape(oriOutputShapeList);
291- conv3dTilingEngine.SetPadding(paddingList);
292- conv3dTilingEngine.SetStride(strideList);
293- conv3dTilingEngine.SetDilation(dilationList);
294- conv3dTilingEngine.SetHF32(enableHF32);
295- if constexpr (std::is_same_v<T, c10::Half>) {
296- conv3dTilingEngine.SetDataType(
297- Conv3dApiTiling::ConvDtype::FLOAT16, Conv3dApiTiling::ConvDtype::FLOAT16,
298- Conv3dApiTiling::ConvDtype::FLOAT16);
299- } else if constexpr (std::is_same_v<T, c10::BFloat16>) {
300- conv3dTilingEngine.SetDataType(
301- Conv3dApiTiling::ConvDtype::BF16, Conv3dApiTiling::ConvDtype::BF16, Conv3dApiTiling::ConvDtype::BF16);
302- } else {
303- conv3dTilingEngine.SetDataType(
304- Conv3dApiTiling::ConvDtype::FLOAT32, Conv3dApiTiling::ConvDtype::FLOAT32,
305- Conv3dApiTiling::ConvDtype::FLOAT32);
306- }
307- if (bias.has_value()) {
308- if constexpr (std::is_same_v<T, c10::BFloat16> || std::is_same_v<T, float>) {
309- bias_ptr = (__gm__ uint8_t*)bias.value().data_ptr<float>();
310- conv3dTilingEngine.SetBias(true, Conv3dApiTiling::ConvDtype::FLOAT32);
311- } else {
312- bias_ptr = (__gm__ uint8_t*)bias.value().data_ptr<T>();
313- conv3dTilingEngine.SetBias(true, Conv3dApiTiling::ConvDtype::FLOAT16);
314- }
315- }
316- Conv3DV2TilingData tilingData;
317- if (!conv3dTilingEngine.GetConv3DV2TilingData(tilingData)) {
318- throw std::runtime_error("Failed to GetConv3DV2TilingData!");
319- }
320- uint32_t blockDim = tilingData.conv3dRunInfo.batchDim * tilingData.conv3dRunInfo.doDim *
321- tilingData.conv3dRunInfo.mDim * tilingData.conv3dRunInfo.nDim;
322- Conv3dCustomKernel<T><<<blockDim, nullptr, stream>>>(input_ptr, weight_ptr, bias_ptr, output_ptr, tilingData);
323-}
324- 
325-template <>
326-void Conv3dCustomApi<double>(
327- aclrtStream stream, const at::Tensor& input, const at::Tensor& weight, const torch::Tensor& output,
328- const vector<int64_t> strideList, const vector<int64_t> paddingList, const vector<int64_t> dilationList,
329- const vector<int64_t> oriInputShapeList, const vector<int64_t> oriWeightShapeList,
330- const c10::optional<torch::Tensor>& bias, bool enableHF32)
331-{
332- throw std::runtime_error("Double is not supported on aicore!");
333-}
334- 
335-torch::Tensor CreateOutputTensor(
336- const torch::Tensor& input, const c10::IntArrayRef& oriInputShape, const c10::IntArrayRef& oriWeightShape,
337- const c10::IntArrayRef& strides, const c10::IntArrayRef& pads, const c10::IntArrayRef& dilations)
338-{
339- int64_t N = oriInputShape[N_DIM_IDX_NCDHW];
340- int64_t D = oriInputShape[D_DIM_IDX_NCDHW];
341- int64_t H = oriInputShape[H_DIM_IDX_NCDHW];
342- int64_t W = oriInputShape[W_DIM_IDX_NCDHW];
343- int64_t Co = oriWeightShape[N_DIM_IDX_NCDHW];
344- vector<int64_t> kernelSize = {
345- oriWeightShape[D_DIM_IDX_NCDHW], oriWeightShape[H_DIM_IDX_NCDHW], oriWeightShape[W_DIM_IDX_NCDHW]};
346- 
347- int64_t Do = (D + 2 * pads[ATTRS_D_DIM_IDX_NCDHW] -
348- dilations[ATTRS_D_DIM_IDX_NCDHW] * (kernelSize[ATTRS_D_DIM_IDX_NCDHW] - 1) - 1) /
349- strides[ATTRS_D_DIM_IDX_NCDHW] +
350- 1;
351- int64_t Ho = (H + 2 * pads[ATTRS_H_DIM_IDX_NCDHW] -
352- dilations[ATTRS_H_DIM_IDX_NCDHW] * (kernelSize[ATTRS_H_DIM_IDX_NCDHW] - 1) - 1) /
353- strides[ATTRS_H_DIM_IDX_NCDHW] +
354- 1;
355- int64_t Wo = (W + 2 * pads[ATTRS_W_DIM_IDX_NCDHW] -
356- dilations[ATTRS_W_DIM_IDX_NCDHW] * (kernelSize[ATTRS_W_DIM_IDX_NCDHW] - 1) - 1) /
357- strides[ATTRS_W_DIM_IDX_NCDHW] +
358- 1;
359- 
360- TORCH_CHECK(N > 0, "Output of N has to be positive, but got ", N);
361- TORCH_CHECK(Co > 0, "Output of C has to be positive, but got ", Co);
362- TORCH_CHECK(Do > 0, "Output of D has to be positive, but got ", Do);
363- TORCH_CHECK(Ho > 0, "Output of H has to be positive, but got ", Ho);
364- TORCH_CHECK(Wo > 0, "Output of W has to be positive, but got ", Wo);
365- 
366- uint8_t k0 = K0_BY_DTYPE_MAP[input.scalar_type()];
367- 
368- return torch::zeros({N, Do, static_cast<int64_t>(Conv3dApiTiling::CeilDiv(Co, k0)), Ho, Wo, k0}, input.options());
369-}
370- 
371-torch::Tensor Conv3dCustomNpu(
372- const torch::Tensor& input, const torch::Tensor& weight, const c10::IntArrayRef& strides,
373- const c10::IntArrayRef& pads, const c10::IntArrayRef& dilations, const c10::IntArrayRef& oriInputShape,
374- const c10::IntArrayRef& oriWeightShape, const c10::optional<torch::Tensor>& bias, bool enableHF32)
375-{
376- TORCH_CHECK(torch_npu::utils::is_npu(input), "Input tensor must be on NPU device");
377- TORCH_CHECK(torch_npu::utils::is_npu(weight), "Weight tensor must be on NPU device");
378- if (bias.has_value()) {
379- TORCH_CHECK(torch_npu::utils::is_npu(bias.value()), "Bias tensor must be on NPU device");
380- if (input.scalar_type() == at::kBFloat16) {
381- TORCH_CHECK(
382- bias.value().scalar_type() == at::kFloat, "Bias type must be float when input dtype is bfloat16");
383- }
384- }
385- TORCH_CHECK(input.dim() == NDC1HWC0_DIM, "Input must be 5D (N, D, C1, H, W, C0)");
386- TORCH_CHECK(weight.dim() == FRACTAL_Z_3D_DIM, "Weight must be 4D (Kd*C1*Kh*Kw, N1, N0, C0)");
387- TORCH_CHECK(strides.size() == CONV_ATTRS_DIM, "Stride must have 3 elements");
388- TORCH_CHECK(pads.size() == CONV_ATTRS_DIM, "Padding must have 3 elements");
389- TORCH_CHECK(dilations.size() == CONV_ATTRS_DIM, "Dilation must have 3 elements");
390- TORCH_CHECK(oriInputShape.size() == NCDHW_DIM, "Origin Input must have 5 elements (N, C, D, H, W)");
391- TORCH_CHECK(oriWeightShape.size() == NCDHW_DIM, "Origin Weight must have 5 elements (N, C, D, H, W)");
392- 
393- auto output = CreateOutputTensor(input, oriInputShape, oriWeightShape, strides, pads, dilations);
394- auto stream = c10_npu::getCurrentNPUStream().stream(false);
395- vector<int64_t> oriInputShapeList(oriInputShape.begin(), oriInputShape.end());
396- vector<int64_t> oriWeightShapeList(oriWeightShape.begin(), oriWeightShape.end());
397- vector<int64_t> strideList(strides.begin(), strides.end());
398- vector<int64_t> paddingList = {pads[ATTRS_D_DIM_IDX_NCDHW], pads[ATTRS_D_DIM_IDX_NCDHW],
399- pads[ATTRS_H_DIM_IDX_NCDHW], pads[ATTRS_H_DIM_IDX_NCDHW],
400- pads[ATTRS_W_DIM_IDX_NCDHW], pads[ATTRS_W_DIM_IDX_NCDHW]};
401- vector<int64_t> dilationList(dilations.begin(), dilations.end());
402- 
403- auto acl_call = [=]() -> int {
404- AT_DISPATCH_FLOATING_TYPES_AND2(at::kHalf, at::kBFloat16, input.scalar_type(), "Conv3dCustomNpu", [&] {
405- Conv3dCustomApi<scalar_t>(
406- stream, input, weight, output, strideList, paddingList, dilationList, oriInputShapeList,
407- oriWeightShapeList, bias, enableHF32);
408- });
409- return 0;
410- };
411- at_npu::native::OpCommand::RunOpApiV2("Conv3dv2", acl_call);
412- return output;
413-}
414- 
415-TORCH_LIBRARY_FRAGMENT(EXTENSION_MODULE_NAME, m)
416-{
417- m.def(
418- "conv3d_custom(Tensor input, Tensor weight, int[3] stride, int[3] padding, int[3] dilation, int[5] "
419- "oriInputShape, int[5] oriWeightShape, Tensor? bias, bool enable_hf32 = False) -> Tensor");
420-}
421- 
422-// Register Ascend implementations for conv3dv2
423-TORCH_LIBRARY_IMPL(EXTENSION_MODULE_NAME, PrivateUse1, m)
424-{
425- m.impl("conv3d_custom", Conv3dCustomNpu);
426-}
427- 
428-} // namespace Conv3dCustom
429-} // namespace ascend_ops
Dexamples/fast_kernel_launch_example/tests/conv3d_custom/test_conv3d_custom.py+0-93
@@ -1,93 +0,0 @@
1-#!/usr/bin/env python3
2-# -*- coding: UTF-8 -*-
3-# ----------------------------------------------------------------------------
4-# This program is free software, you can redistribute it and/or modify it.
5-# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6-# This file is a part of the CANN Open Software.
7-# Licensed under CANN Open Software License Agreement Version 2.0 (the "License").
8-# Please refer to the License for details. You may not use this file except in compliance with the License.
9-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING
10-# BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
11-# See LICENSE in the root of the software repository for the full text of the License.
12-# ----------------------------------------------------------------------------
13- 
14-import torch
15-import torch_npu
16-import ascend_ops
17- 
18- 
19-def cpu_conv3d(input_tensor, weight_tensor, bias, fmap_shape, weight_shape, bias_flag, stride,
20- padding, dilation, groups, dtype):
21- cpu_dtype = dtype
22- if dtype == torch.float16:
23- cpu_dtype = torch.float32
24- input_tensor = input_tensor.to(dtype=cpu_dtype)
25- weight_tensor = weight_tensor.to(dtype=cpu_dtype)
26- bias = bias.to(dtype=cpu_dtype)
27- with torch.no_grad():
28- model0 = torch.nn.Conv3d(fmap_shape[1], weight_shape[0],
29- (weight_shape[2], weight_shape[3], weight_shape[4]),
30- bias=bias_flag, stride=stride, padding=padding, dilation=dilation,
31- groups=groups, dtype=cpu_dtype)
32- model0.weight.data = weight_tensor
33- if bias_flag:
34- model0.bias.data = bias
35- output_tensor = model0(input_tensor)
36- if dtype == torch.float16:
37- output_tensor = output_tensor.to(torch.float16)
38- return output_tensor
39- 
40- 
41-def test_conv3d_custom():
42- for data_type in [torch.float16, torch.bfloat16, torch.float32]:
43- input_tensor = torch.randn(1, 1, 16, 16, 16).to(data_type)
44- weight_tensor = torch.randn(1, 1, 3, 3, 3).to(data_type)
45- bias_tensor = torch.randn(1).to(data_type)
46- stride = [1, 1, 1]
47- dilation = [1, 1, 1]
48- padding = [0, 0, 0]
49- 
50- print(f"Dtype: {data_type}")
51- print(f"Input shape: {input_tensor.shape}")
52- print(f"Weight shape: {weight_tensor.shape}")
53- print(f"Bias shape: {bias_tensor.shape}")
54- 
55- cpu_result = cpu_conv3d(input_tensor, weight_tensor, bias_tensor, input_tensor.shape, weight_tensor.shape, True,
56- stride, padding, dilation, 1, data_type)
57- if input_tensor.dtype == torch.bfloat16 and bias_tensor:
58- bias_tensor = bias_tensor.to(torch.float32) # bflost16场景的bias仅支持float32类型
59- npu_result = ascend_ops.ops.conv3d_custom(input_tensor, weight_tensor, stride, padding, dilation, bias_tensor)
60- print(f"CPU Result shape: {cpu_result.shape}")
61- print(f"NPU Result shape: {npu_result.shape}")
62- print(f"compare CPU Result vs NPU Result: "
63- f"{torch.allclose(cpu_result, npu_result, rtol=1e-02, atol=1e-02, equal_nan=True)}\n")
64- 
65- 
66-def test_conv3d_custom_by_hf32():
67- for data_type in [torch.float32]:
68- input_tensor = torch.randn(1, 1, 16, 16, 16).to(data_type)
69- weight_tensor = torch.randn(1, 1, 3, 3, 3).to(data_type)
70- bias_tensor = torch.randn(1).to(data_type)
71- stride = [1, 1, 1]
72- dilation = [1, 1, 1]
73- padding = [0, 0, 0]
74- 
75- print(f"Dtype: {data_type}")
76- print(f"Input shape: {input_tensor.shape}")
77- print(f"Weight shape: {weight_tensor.shape}")
78- print(f"Bias shape: {bias_tensor.shape}")
79- 
80- cpu_result = cpu_conv3d(input_tensor, weight_tensor, bias_tensor, input_tensor.shape, weight_tensor.shape, True,
81- stride, padding, dilation, 1, data_type)
82- enable_hf32 = True
83- npu_result = ascend_ops.ops.conv3d_custom(input_tensor, weight_tensor, stride, padding, dilation, bias_tensor,
84- enable_hf32)
85- print(f"CPU Result shape: {cpu_result.shape}")
86- print(f"NPU Result shape: {npu_result.shape}")
87- print(f"compare CPU Result vs NPU Result: "
88- f"{torch.allclose(cpu_result, npu_result, rtol=1e-02, atol=1e-02, equal_nan=True)}\n")
89- 
90- 
91-if __name__ == '__main__':
92- test_conv3d_custom()
93- test_conv3d_custom_by_hf32()