This change adds English version of the documentation-triton-ascend-AtomGit

ACONTRIBUTING.md+168-0

@@ -0,0 +1,168 @@
	1	+# Triton Ascend Contribution Guide
	2	+
	3	+- [Getting Started](#getting-started.md)
	4	+- [Developer Guide](#developer-guide.md)
	5	+ - [Coding Style](#coding-style.md)
	6	+ - [Fork-Pull Mode](#forkpull-mode.md)
	7	+ - [Troubleshooting Gated Commit](#troubleshooting-gated-commit.md)
	8	+ - [Issue Specifications](#issue-specifications.md)
	9	+ - [Pull Request Proposal](#pull-request-proposal.md)
	10	+
	11	+<h2 id="getting-started.md">Getting Started</h2>
	12	+
	13	+- Fork the Triton Ascend repository on [GitCode](https://gitcode.com/Ascend/triton-ascend).
	14	+- Check the [README.md](https://gitcode.com/Ascend/triton-ascend/blob/master/README.md) file to obtain the project information and build the development environment.
	15	+
	16	+
	17	+
	18	+<h2 id="developer-guide.md">Developer Guide</h2>
	19	+
	20	+- [Coding Style](#coding-style.md)
	21	+- [Fork-Pull Mode](#forkpull-mode.md)
	22	+- [Troubleshooting Gated Commit](#troubleshooting-gated-commit.md)
	23	+- [Issue Specifications](#issue-specifications.md)
	24	+- [Pull Request Proposal](#pull-request-proposal.md)
	25	+
	26	+
	27	+
	28	+<h2 id="coding-style .md">Coding Style</h2>
	29	+
	30	+Follow the coding style below to make Triton Ascend easy to develop, maintain, and review.
	31	+
	32	+- Coding Guide
	33	+
	34	+ Use the unified coding style of the Triton Ascend community. The recommended coding style for Python is [PEP 8 Coding Style](https://pep8.org/), and that for C++ is [LLVM Coding Standards](https://llvm.org/docs/CodingStandards.html). You can use [clang-tidy](https://github.com/llvm/llvm-project/blob/main/.clang-tidy), [CppLint](https://github.com/cpplint/cpplint), [CppCheck](http://cppcheck.sourceforge.net/), [CMakeLint](https://github.com/cmake-lint/cmake-lint), [CodeSpell](https://github.com/codespell-project/codespell), [ShellCheck](https://github.com/koalaman/shellcheck), and [pylint](https://pylint.org/) to check the code format. You are advised to install these plug-ins in your IDE.
	35	+
	36	+- Unit Test Guide
	37	+
	38	+ Use the unified unit test style of the Triton Ascend community. The recommended unit test style for Python is [pytest](http://www.pytest.org/en/latest/), and that for C++ is [GoogleTest Primer](https://github.com/google/googletest/blob/master/docs/primer.md). The design intent of a test case should be reflected by its annotation name. For details about how to design test cases, see [gather Test Cases](https://gitcode.com/Ascend/triton-ascend/blob/master/ascend/examples/pytest_ut/test_gather.py) and [layer_norm Test Cases](https://gitcode.com/Ascend/triton-ascend/blob/master/ascend/examples/tutorials/03-layer-norm.py).
	39	+
	40	+- Refactoring Guide
	41	+
	42	+ We encourage developers to refactor our code to eliminate code smells. The refactored code should also comply with the coding style and testing style requirements. When receiving a warning, refactor the code to be merged.
	43	+
	44	+
	45	+
	46	+<h2 id="Fork-Pull Mode .md">Fork-Pull Mode</h2>
	47	+
	48	+1. Fork the Triton Ascend project.
	49	+
	50	+Before committing your code to the Triton Ascend project, ensure that you have forked the Triton Ascend project to your own repository. You will then develop the project in your own forked repository and merge it to the Triton Ascend project through a pull request (PR). This means that there is parallel development between the Triton Ascend repository and your own repository. Be careful to avoid inconsistency between repositories.
	51	+
	52	+2. Clone a remote repository.
	53	+
	54	+Use git to clone the Triton Ascend project you have forked and add the upstream repository.
	55	+
	56	+```shell
	57	+git clone https://github.com/triton-lang/{your_forked_repo}/triton-ascend.git && cd triton-ascend && git submodule update --init --depth 1
	58	+git remote add upstream https://github.com/triton-lang/triton-ascend.git
	59	+```
	60	+
	61	+3. Develop code locally.
	62	+
	63	+Before developing your code, you need to set up the development environment according to the [Triton Ascend Installation Guide](https://gitcode.com/Ascend/triton-ascend/blob/main/docs/en/installation_guide.md).
	64	+
	65	+To avoid inconsistency between branches, create a new local development branch for new features.
	66	+
	67	+```shell
	68	+git checkout -b {new_branch_name} origin/master
	69	+git fetch upstream #Fetch the latest code from the upstream repository
	70	+git rebase upstream/master #Rebase onto the latest upstream
	71	+```
	72	+
	73	+Taking the master branch as an example, Triton Ascend may create version branches or downstream development branches as required. After creating a branch and synchronizing the upstream master branch, you can start developing your code.
	74	+
	75	+4. Perform a self-test.
	76	+
	77	+After the code is modified, check whether the changes can pass the test.
	78	+
	79	+Write a test script for the developed code in the ascend/examples/pytest_ut directory of your local code branch, and verify the test script in the local environment to ensure that the changes can pass the test.
	80	+
	81	+5. Push code to the remote repository.
	82	+
	83	+After updating and testing the code, push your commit to the remote repository.
	84	+
	85	+```shell
	86	+git add .
	87	+git status #Check the updated files
	88	+git commit -m "Your commit title"
	89	+git commit -s --amend #Add the concrete description of your commit
	90	+git push origin {your_new_branch_name}
	91	+```
	92	+
	93	+6. Create a pull request to the Triton Ascend main repository.
	94	+
	95	+After pushing code to your remote repository, create a pull request between your new branch and the Triton Ascend master branch. After the merge request is created, Jenkins CI will be automatically set to build your pipeline test. You are advised to merge your pull request to the upstream master branch as soon as possible to reduce the merge risk.
	96	+
	97	+The pipeline execution process after a PR is committed is as follows:
	98	+
	99	+- Comment /compile to start the pipeline test. If the test fails, modify the code as prompted and comment /compile again to trigger the pipeline test. After the test is passed, the tag ci-pipeline-passed is added.
	100	+
	101	+ ```shell
	102	+ /compile
	103	+ ```
	104	+
	105	+- If SC-FAIL is displayed, check the modification and comment compile#openlibing to manually trigger the check. After the check is passed, the tag SC-SUCC is added.
	106	+
	107	+ ```tex
	108	+ compile#openlibing
	109	+ ```
	110	+
	111	+- After the pipeline passes the test (the ci-pipeline-passed and SC-SUCC tags are added), comment @committers as prompted to review the code so that the code can be quickly merged.
	112	+
	113	+
	114	+
	115	+<h2 id="Troubleshooting Gated Commit.md">Troubleshooting Gated Commit</h2>
	116	+
	117	+Gated commit may encounter the following exceptions. Rectify the exceptions according to the related information.
	118	+
	119	+- Compilation failed
	120	+
	121	+ Check the cause of the compilation failure as prompted, and then recompile the code.
	122	+
	123	+- Static check failed
	124	+
	125	+ Find and fix the exception information in the code as prompted.
	126	+
	127	+- CI pipeline failed
	128	+
	129	+ Find the failed test cases of the CI pipeline as prompted, and then check the cause. After the fault is rectified, run the CI pipeline again.
	130	+
	131	+
	132	+
	133	+<h2 id="issue-specifications.md">Issue Specifications</h2>
	134	+
	135	+A good way to contribute to the project is to send a detailed report when you encounter a problem. We are always very grateful for detailed and thorough bug reports, and we will be very grateful to you for that!
	136	+
	137	+Please include the following information when you file an issue:
	138	+
	139	+- What is the software version (Triton Ascend, Python, OS, etc.) used in your environment?
	140	+- Is it a bug report or a functional request?
	141	+- What kind of issue are you reporting? Add the corresponding tag to highlight it on the issue dashboard.
	142	+- What's happened?
	143	+- What did you expect to happen?
	144	+- How to reproduce the issue? (As accurately as possible)
	145	+
	146	+For details about the templates for filling in issues of different categories, see [Issue Specifications](https://gitcode.com/Ascend/triton-ascend/issues/create/choose).
	147	+
	148	+Notes for contributors:
	149	+
	150	+- If you find an unresolved issue that is exactly what you are trying to solve, comment on the issue and tell others that you will be responsible for handling it.
	151	+- If the issue has existed for a period of time, it is recommended that you perform a pre-check before solving the issue.
	152	+- If you have resolved the issue you report, inform others before closing the issue.
	153	+
	154	+
	155	+<h2 id="pull-request-proposal.md">Pull Request Proposal</h2>
	156	+
	157	+- Propose your ideas as issues.
	158	+- If the new feature to be developed requires a large number of design details, you should also commit the design solution.
	159	+- After reaching a consensus in the issue discussion and design solution review, fork the project and commit a PR.
	160	+- No PR is allowed until you receive 2+LGTM (Looks Good To Me) from the approver. Note that you are not allowed to add LGTM to your own PRs.
	161	+- After the PR is fully discussed, it will be merged, rejected, or abandoned based on the discussion result.
	162	+
	163	+### Notes:
	164	+
	165	+- Avoid any irrelevant changes.
	166	+- Ensure that your commit history is concise and orderly.
	167	+- Before creating a PR, please rebase the latest code from the upstream repository.
	168	+- For a bug-fixing PR, ensure that all related issues and PRs are linked.

RCONTRIBUTING.zh.md→CONTRIBUTING_zh.md+189-189

	@@ -1,189 +1,189 @@
1	-# Triton Ascend贡献指南	1	+# Triton Ascend贡献指南
2	-	2	+
3	-- [贡献者许可协议](#贡献者许可协议.md)	3	+- [贡献者许可协议](#贡献者许可协议.md)
4	-- [入门](#入门.md)	4	+- [入门](#入门.md)
5	-- [开发指导](#开发指导.md)	5	+- [开发指导](#开发指导.md)
6	- - [代码风格](#代码风格.md)	6	+ - [代码风格](#代码风格.md)
7	- - [Fork-Pull开发模式](#Fork-Pull开发模式.md)	7	+ - [Fork-Pull开发模式](#Fork-Pull开发模式.md)
8	- - [代码门禁异常处理](#代码门禁异常处理.md)	8	+ - [代码门禁异常处理](#代码门禁异常处理.md)
9	- - [ISSUE规范](#ISSUE规范.md)	9	+ - [ISSUE规范](#ISSUE规范.md)
10	- - [提出PR](#提出PR.md)	10	+ - [提出PR](#提出PR.md)
11	-	11	+
12	-<h2 id="贡献者许可协议.md">	12	+<h2 id="贡献者许可协议.md">
13	- 贡献者许可协议	13	+ 贡献者许可协议
14	-</h2>	14	+</h2>
15	-	15	+
16	-在您第一次向Triton Ascend社区提交代码之前，需要签署CLA。	16	+在您第一次向Triton Ascend社区提交代码之前，需要签署CLA。
17	-	17	+
18	-对于个人贡献者，签署CLA详细信息参考 [cla使用指南](https://gitcode.com/Ascend/infrastructure/blob/master/docs/cla/cla使用指南.md#faq)	18	+对于个人贡献者，签署CLA详细信息参考 [cla使用指南](https://gitcode.com/Ascend/infrastructure/blob/master/docs/cla/cla使用指南.md#faq)
19	-	19	+
20	-CLA签署地址 [sign](https://clasign.osinfra.cn/sign/690ca9ddf91c03dee6082ab1)	20	+CLA签署地址 [sign](https://clasign.osinfra.cn/sign/690ca9ddf91c03dee6082ab1)
21	-	21	+
22	-	22	+
23	-	23	+
24	-<h2 id="入门.md">入门</h2>	24	+<h2 id="入门.md">入门</h2>
25	-	25	+
26	-- 在[GitCode](https://gitcode.com/Ascend/triton-ascend)上Fork Triton Ascend存储库。	26	+- 在[GitCode](https://gitcode.com/Ascend/triton-ascend)上Fork Triton Ascend存储库。
27	-- 阅读[README.md](https://gitcode.com/Ascend/triton-ascend/blob/master/README.md)获取项目信息和构建开发环境。	27	+- 阅读[README.md](https://gitcode.com/Ascend/triton-ascend/blob/master/README.md)获取项目信息和构建开发环境。
28	-	28	+
29	-	29	+
30	-	30	+
31	-<h2 id="开发指导.md">开发指导</h2>	31	+<h2 id="开发指导.md">开发指导</h2>
32	-	32	+
33	-- [代码风格](#代码风格.md)	33	+- [代码风格](#代码风格.md)
34	-- [Fork-Pull开发模式](#Fork-Pull开发模式.md)	34	+- [Fork-Pull开发模式](#Fork-Pull开发模式.md)
35	-- [代码门禁异常处理](#代码门禁异常处理.md)	35	+- [代码门禁异常处理](#代码门禁异常处理.md)
36	-- [ISSUE规范](#ISSUE规范.md)	36	+- [ISSUE规范](#ISSUE规范.md)
37	-- [提出PR](#提出PR.md)	37	+- [提出PR](#提出PR.md)
38	-	38	+
39	-	39	+
40	-	40	+
41	-<h2 id="代码风格.md">代码风格</h2>	41	+<h2 id="代码风格.md">代码风格</h2>
42	-	42	+
43	-请遵循以下编码风格，以使得Triton Ascend易于开发、维护和审查。	43	+请遵循以下编码风格，以使得Triton Ascend易于开发、维护和审查。
44	-	44	+
45	-- 编码指南	45	+- 编码指南
46	-	46	+
47	- 请使用Triton Ascend社区统一的编码风格，python建议的编码风格是[PEP 8编码样式](https://pep8.org/)，C++编码所建议的风格是 [LLVM 编码规范](https://llvm.org/docs/CodingStandards.html)) 。可以使用[clang-tidy](https://github.com/llvm/llvm-project/blob/main/.clang-tidy)，[CppLint](https://github.com/cpplint/cpplint)，[CppCheck](http://cppcheck.sourceforge.net/)，[CMakeLint](https://github.com/cmake-lint/cmake-lint)，[CodeSpell](https://github.com/codespell-project/codespell)，[ShellCheck](https://github.com/koalaman/shellcheck)和[pylint](https://pylint.org/)检查代码的格式，建议在您的IDE中安装这些插件。	47	+ 请使用Triton Ascend社区统一的编码风格，python建议的编码风格是[PEP 8编码样式](https://pep8.org/)，C++编码所建议的风格是 [LLVM 编码规范](https://llvm.org/docs/CodingStandards.html)) 。可以使用[clang-tidy](https://github.com/llvm/llvm-project/blob/main/.clang-tidy)，[CppLint](https://github.com/cpplint/cpplint)，[CppCheck](http://cppcheck.sourceforge.net/)，[CMakeLint](https://github.com/cmake-lint/cmake-lint)，[CodeSpell](https://github.com/codespell-project/codespell)，[ShellCheck](https://github.com/koalaman/shellcheck)和[pylint](https://pylint.org/)检查代码的格式，建议在您的IDE中安装这些插件。
48	-	48	+
49	-- 单元测试指南	49	+- 单元测试指南
50	-	50	+
51	- 请使用Triton Ascend社区统一的单元测试风格，python建议的单元测试风格是[pytest](http://www.pytest.org/en/latest/)，C++建议的单元测试风格是[Googletest Primer](#https://github.com/google/googletest/blob/master/docs/primer.md)。测试用例的设计意图应该通过它的注释名称来反映。测试用例的设计请参考[gather测试用例](https://gitcode.com/Ascend/triton-ascend/blob/master/ascend/examples/pytest_ut/test_gather.py)，[layer_norm测试用例](https://gitcode.com/Ascend/triton-ascend/blob/master/ascend/examples/tutorials/03-layer-norm.py)	51	+ 请使用Triton Ascend社区统一的单元测试风格，python建议的单元测试风格是[pytest](http://www.pytest.org/en/latest/)，C++建议的单元测试风格是[Googletest Primer](#https://github.com/google/googletest/blob/master/docs/primer.md)。测试用例的设计意图应该通过它的注释名称来反映。测试用例的设计请参考[gather测试用例](https://gitcode.com/Ascend/triton-ascend/blob/master/ascend/examples/pytest_ut/test_gather.py)，[layer_norm测试用例](https://gitcode.com/Ascend/triton-ascend/blob/master/ascend/examples/tutorials/03-layer-norm.py)
52	-	52	+
53	-- 重构指南	53	+- 重构指南
54	-	54	+
55	- 我们鼓励开发人员对我们的代码进行重构来消除【代码坏味道】。重构的代码也应该遵循编码风格和测试风格的要求。当您收到警告时，您需要重构要合并的代码。	55	+ 我们鼓励开发人员对我们的代码进行重构来消除【代码坏味道】。重构的代码也应该遵循编码风格和测试风格的要求。当您收到警告时，您需要重构要合并的代码。
56	-	56	+
57	-	57	+
58	-	58	+
59	-<h2 id="Fork-Pull开发模式.md">Fork-Pull开发模式</h2>	59	+<h2 id="Fork-Pull开发模式.md">Fork-Pull开发模式</h2>
60	-	60	+
61	-1、Fork Triton Ascend项目	61	+1、Fork Triton Ascend项目
62	-	62	+
63	-在您向Triton Ascend项目提交自己的代码之前，请确保已经将Triton Ascend项目Fork到您自己的存储库。后续您将在自己Fork的项目上进行开发，并通过Pull Request的方式合并到Triton Ascend项目。这意味着Triton Ascend存储库和您自己的存储库之间存在并行开发，因此请注意保持存储库之间的一致性。	63	+在您向Triton Ascend项目提交自己的代码之前，请确保已经将Triton Ascend项目Fork到您自己的存储库。后续您将在自己Fork的项目上进行开发，并通过Pull Request的方式合并到Triton Ascend项目。这意味着Triton Ascend存储库和您自己的存储库之间存在并行开发，因此请注意保持存储库之间的一致性。
64	-	64	+
65	-2、克隆远程仓库	65	+2、克隆远程仓库
66	-	66	+
67	-使用git克隆您fork的Triton Ascend项目&添加上游仓库upstream：	67	+使用git克隆您fork的Triton Ascend项目&添加上游仓库upstream：
68	-	68	+
69	-```shell	69	+```shell
70	-git clone https://gitcode.com/{your_forked_repo}/triton-ascend.git && cd triton-ascend && git submodule update --init --depth 1	70	+git clone https://gitcode.com/{your_forked_repo}/triton-ascend.git && cd triton-ascend && git submodule update --init --depth 1
71	-git remote add upstream https://gitcode.com/Ascend/triton-ascend.git	71	+git remote add upstream https://gitcode.com/Ascend/triton-ascend.git
72	-```	72	+```
73	-	73	+
74	-3、本地环境开发代码	74	+3、本地环境开发代码
75	-	75	+
76	-在开发您的代码之前，您需要根据[Triton Ascend安装指南](https://gitcode.com/Ascend/triton-ascend/blob/main/docs/zh/installation_guide.md)搭建开发环境。	76	+在开发您的代码之前，您需要根据[Triton Ascend安装指南](https://gitcode.com/Ascend/triton-ascend/blob/main/docs/zh/installation_guide.md)搭建开发环境。
77	-	77	+
78	-为避免多个分支间的不一致问题，请创建新的本地开发分支进行新特性的开发：	78	+为避免多个分支间的不一致问题，请创建新的本地开发分支进行新特性的开发：
79	-	79	+
80	-```shell	80	+```shell
81	-git checkout -b {new_branch_name} origin/master	81	+git checkout -b {new_branch_name} origin/master
82	-git fetch upstream #Fetch the latest code from the upstream repository	82	+git fetch upstream #Fetch the latest code from the upstream repository
83	-git rebase upstream/master #Rebase onto the latest upstream	83	+git rebase upstream/master #Rebase onto the latest upstream
84	-```	84	+```
85	-	85	+
86	-以master分支为例，Triton Ascend可能会根据需要创建版本分支或下游开发分支。当您创建完分支&同步上游master分支更新后，就可以开始开发您的代码了。	86	+以master分支为例，Triton Ascend可能会根据需要创建版本分支或下游开发分支。当您创建完分支&同步上游master分支更新后，就可以开始开发您的代码了。
87	-	87	+
88	-4、代码更改自测	88	+4、代码更改自测
89	-	89	+
90	-完成代码更改后，请检查您的更改是否可以通过测试：	90	+完成代码更改后，请检查您的更改是否可以通过测试：
91	-	91	+
92	-在本地代码分支的ascend/examples/pytest_ut路径下为您开发的代码编写测试用例代码，并在本地环境中验证您的测试脚本，确保您的更改可以通过测试。	92	+在本地代码分支的ascend/examples/pytest_ut路径下为您开发的代码编写测试用例代码，并在本地环境中验证您的测试脚本，确保您的更改可以通过测试。
93	-	93	+
94	-5、代码推送到远程仓库	94	+5、代码推送到远程仓库
95	-	95	+
96	-代码更新&测试完成后，推送您的commit到您的远程仓库。	96	+代码更新&测试完成后，推送您的commit到您的远程仓库。
97	-	97	+
98	-```shell	98	+```shell
99	-git add .	99	+git add .
100	-git status #Check the updated files	100	+git status #Check the updated files
101	-git commit -m "Your commit title"	101	+git commit -m "Your commit title"
102	-git commit -s --amend #Add the concrete description of your commit	102	+git commit -s --amend #Add the concrete description of your commit
103	-git push origin {your_new_branch_name}	103	+git push origin {your_new_branch_name}
104	-```	104	+```
105	-	105	+
106	-6、向Triton Ascend主仓创建拉取请求	106	+6、向Triton Ascend主仓创建拉取请求
107	-	107	+
108	-代码推送至您的远程仓库后，您需要在您的新分支和Triton Ascend master分支之间新建Pull Request。完成新建合并请求后，“Jenkins CI“将自动设置为您构建流水线测试。您的Pull Request请尽快合并到上游master分支，以降低合并风险。	108	+代码推送至您的远程仓库后，您需要在您的新分支和Triton Ascend master分支之间新建Pull Request。完成新建合并请求后，“Jenkins CI“将自动设置为您构建流水线测试。您的Pull Request请尽快合并到上游master分支，以降低合并风险。
109	-	109	+
110	-提交PR后流水线执行命令流程	110	+提交PR后流水线执行命令流程
111	-	111	+
112	-- 如果PR的标签显示ascend-cla/no，签署cla后评论/check-cla检查cla签署状态，cla签署成功后获得标签 ascend-cla/yes。	112	+- 如果PR的标签显示ascend-cla/no，签署cla后评论/check-cla检查cla签署状态，cla签署成功后获得标签 ascend-cla/yes。
113	-	113	+
114	- ```shell	114	+ ```shell
115	- /check-cla	115	+ /check-cla
116	- ```	116	+ ```
117	-	117	+
118	-- 评论/compile启动流水线测试，如果未通过测试，根据提示修改后再次评论/compile触发流水线测试，通过后获得标签 ci-pipeline-passed。	118	+- 评论/compile启动流水线测试，如果未通过测试，根据提示修改后再次评论/compile触发流水线测试，通过后获得标签 ci-pipeline-passed。
119	-	119	+
120	- ```shell	120	+ ```shell
121	- /compile	121	+ /compile
122	- ```	122	+ ```
123	-	123	+
124	-- 如果SC-FAIL，检查修改后可评论compile#openlibing手动触发检查，检查通过后获得标签 SC-SUCC。	124	+- 如果SC-FAIL，检查修改后可评论compile#openlibing手动触发检查，检查通过后获得标签 SC-SUCC。
125	-	125	+
126	- ```tex	126	+ ```tex
127	- compile#openlibing	127	+ compile#openlibing
128	- ```	128	+ ```
129	-	129	+
130	-- 流水线pass之后（收到ci-pipeline-passed、ascend-cla/yes、SC-SUCC标签），根据提示@committers进行代码review，以便快速合入。	130	+- 流水线pass之后（收到ci-pipeline-passed、ascend-cla/yes、SC-SUCC标签），根据提示@committers进行代码review，以便快速合入。
131	-	131	+
132	-	132	+
133	-	133	+
134	-<h2 id="代码门禁异常处理.md">代码门禁异常处理</h2>	134	+<h2 id="代码门禁异常处理.md">代码门禁异常处理</h2>
135	-	135	+
136	-代码门禁异常主要包含以下几种情况，请根据相关提示信息解决门禁异常问题。	136	+代码门禁异常主要包含以下几种情况，请根据相关提示信息解决门禁异常问题。
137	-	137	+
138	-- 编译失败	138	+- 编译失败
139	-	139	+
140	- 请根据提示信息，检查编译失败的原因，解决后重新编译即可。	140	+ 请根据提示信息，检查编译失败的原因，解决后重新编译即可。
141	-	141	+
142	-- 静态检查失败	142	+- 静态检查失败
143	-	143	+
144	- 请根据提示信息，查找出代码中的异常信息并解决。	144	+ 请根据提示信息，查找出代码中的异常信息并解决。
145	-	145	+
146	-- CI流水线未通过	146	+- CI流水线未通过
147	-	147	+
148	- 请根据提示信息，查找出CI流水线未通过的测试用例并检查原因，解决后重新运行CI流水线。	148	+ 请根据提示信息，查找出CI流水线未通过的测试用例并检查原因，解决后重新运行CI流水线。
149	-	149	+
150	-	150	+
151	-	151	+
152	-<h2 id="ISSUE规范.md">ISSUE规范</h2>	152	+<h2 id="ISSUE规范.md">ISSUE规范</h2>
153	-	153	+
154	-为项目做贡献的一个好的方法是在遇到问题时发送详细报告。我们总是非常感谢写得详细、彻底的错误报告，并会因此非常感谢您！	154	+为项目做贡献的一个好的方法是在遇到问题时发送详细报告。我们总是非常感谢写得详细、彻底的错误报告，并会因此非常感谢您！
155	-	155	+
156	-在报告问题时，请参考以下格式：	156	+在报告问题时，请参考以下格式：
157	-	157	+
158	-- 您环境里使用的软件版本（Triton Ascend、python、os等）？	158	+- 您环境里使用的软件版本（Triton Ascend、python、os等）？
159	-- 这是一个错误报告还是功能请求？	159	+- 这是一个错误报告还是功能请求？
160	-- 您报告的是什么样的问题，添加对应的标签以便在问题仪表盘上突出显示？	160	+- 您报告的是什么样的问题，添加对应的标签以便在问题仪表盘上突出显示？
161	-- 发生了什么？	161	+- 发生了什么？
162	-- 您预计会发生什么？	162	+- 您预计会发生什么？
163	-- 如何重现它？（尽可能精确）	163	+- 如何重现它？（尽可能精确）
164	-	164	+
165	-不同类别的ISSUE填写模板请参考[ISSUE填写规范](https://gitcode.com/Ascend/triton-ascend/issues/create/choose)	165	+不同类别的ISSUE填写模板请参考[ISSUE填写规范](https://gitcode.com/Ascend/triton-ascend/issues/create/choose)
166	-	166	+
167	-问题咨询：	167	+问题咨询：
168	-	168	+
169	-- 如果您发现一个未解决的问题，而这个问题正是您要解决的，请对该问题发表评论，告诉其他人您将负责这个问题。	169	+- 如果您发现一个未解决的问题，而这个问题正是您要解决的，请对该问题发表评论，告诉其他人您将负责这个问题。
170	-- 如果问题已经打开一段时间，请您在解决该问题前进行预检查。	170	+- 如果问题已经打开一段时间，请您在解决该问题前进行预检查。
171	-- 如果您解决了自己报告的问题，在关闭该问题前还需要让其他人知道。	171	+- 如果您解决了自己报告的问题，在关闭该问题前还需要让其他人知道。
172	-	172	+
173	-	173	+
174	-	174	+
175	-<h2 id="提出PR.md">提出PR</h2>	175	+<h2 id="提出PR.md">提出PR</h2>
176	-	176	+
177	-- 在[GitCode](https://gitcode.com/Ascend/triton-ascend)上提出您的想法作为问题。	177	+- 在[GitCode](https://gitcode.com/Ascend/triton-ascend)上提出您的想法作为问题。
178	-- 如果要开发的新功能需要大量设计细节，您还应提交设计方案。	178	+- 如果要开发的新功能需要大量设计细节，您还应提交设计方案。
179	-- 在问题讨论和设计方案审查达成共识后，再进行fork开发并提交PR。	179	+- 在问题讨论和设计方案审查达成共识后，再进行fork开发并提交PR。
180	-- 在从Approver那里收到2+LGTM（Looks Good To Me）前不允许任何PR 。请注意审批人不允许在自己的PR上添加LGTM。	180	+- 在从Approver那里收到2+LGTM（Looks Good To Me）前不允许任何PR 。请注意审批人不允许在自己的PR上添加LGTM。
181	-- 在PR被充分讨论后，将根据讨论结果对PR进行合并、拒绝或放弃。	181	+- 在PR被充分讨论后，将根据讨论结果对PR进行合并、拒绝或放弃。
182	-- PR样例：[PR样例](https://gitcode.com/Ascend/triton-ascend/pull/936)	182	+- PR样例：[PR样例](https://gitcode.com/Ascend/triton-ascend/pull/936)
183	-	183	+
184	-### 注意事项：	184	+### 注意事项：
185	-	185	+
186	-- 应避免任何不相关的更改。	186	+- 应避免任何不相关的更改。
187	-- 确保您的提交历史是简洁有序的。	187	+- 确保您的提交历史是简洁有序的。
188	-- 创建PR前请rebase上游仓库最新代码。	188	+- 创建PR前请rebase上游仓库最新代码。
189	-- 对于错误修复 PR，请确保链接所有相关Issue 和 PR。	189	+- 对于错误修复 PR，请确保链接所有相关Issue 和 PR。

MREADME.md+71-71

	@@ -1,117 +1,117 @@
1	# Triton-Ascend	1	# Triton-Ascend
2		2
3	-## 项目简介与价值主张	3	+## Project Overview and Value Proposition
4	-Triton-Ascend是面向昇腾平台构建的Triton编译框架，旨在让Triton代码能够在昇腾硬件上高效运行。	4	+Triton-Ascend is a Triton compilation framework built for the Ascend platform, aiming to enable Triton code to run efficiently on Ascend hardware.
5	-- #### 核心价值说明	5	+- #### Core Value
6	-Triton是近几年来受到开发者青睐的Python化编译框架。开发者仅需关注Tile/Block的切分方式以及基于Tile/Block的运算逻辑，编译器将在Triton代码的编译过程中结合底层硬件特点自动完成内存分配、数据搬运、数据计算、流水并行等，因此，算子的开发难度大幅降低、开发效率显著提升。	6	+Triton is a Python-based compilation framework that has been favored by developers in recent years. Developers only need to focus on the tile/block slicing mode and the computation logic based on tiles/blocks. During the compilation of Triton code, the compiler automatically completes memory allocation, data transfer, data computation, and pipeline parallelism based on the characteristics of underlying hardware. This greatly reduces the operator development difficulty and significantly improves the development efficiency.
7	-Triton-Ascend将Triton编译栈适配到华为昇腾NPU上，在Triton的基础上提供一系列针对性的优化，使Triton代码能够编译后在昇腾硬件上高效运行。	7	+Triton-Ascend adapts the Triton compilation stack to Huawei Ascend NPUs and provides a series of optimizations based on Triton, so that Triton code can run efficiently on Ascend hardware after compilation.
8	-目前，Triton-Ascend仍在持续完善中，我们将不断提升Triton Python API完备度、数据类型支持度、访存方式灵活性等，并持续优化编译器的自动优化能力，提升Triton-Ascend整体的功能与性能泛化性。	8	+Currently, Triton-Ascend is still being improved. We will continuously improve the completeness of Triton Python APIs, support more data types, make memory access more flexible, and continuously optimize the automatic optimization capability of the compiler to improve the overall functionality and performance generalization of Triton-Ascend.
9	-- #### 昇腾生态定位	9	+- #### Ascend Ecosystem Positioning
10	-Triton-Ascend编译框架打通了Triton与昇腾硬件之间的壁垒，使熟悉Triton框架的开发者可以更有效率地使用昇腾NPU。它通过提供通用、高效的算子开发范式，为昇腾软件栈补齐了敏捷开发的关键一环，极大丰富了昇腾的算子库和上层应用生态。	10	+The Triton-Ascend compilation framework removes the barriers between Triton and Ascend hardware, enabling developers who are familiar with the Triton framework to use Ascend NPUs more efficiently. It provides a universal and efficient operator development paradigm, which is a key part of agile development for the Ascend software stack. This greatly enriches the Ascend operator library and upper-layer application ecosystem.
11		11
12	-## 最新动态与里程碑	12	+## Latest Updates and Milestones
13	-- #### 近期版本更新	13	+- #### Latest Updates
14	-当前版本：[Triton-Ascend 3.2.0](https://pypi.org/project/triton-ascend/)	14	+Current version: [Triton-Ascend 3.2.0](https://pypi.org/project/triton-ascend/)
15	-配套CANN版本：[昇腾CANN社区版8.5.0](https://www.hiascend.com/developer/download/community/result?module=cann&cann=8.5.0)	15	+CANN version: [Ascend CANN Community Edition 8.5.0](https://www.hiascend.com/developer/download/community/result?module=cann&cann=8.5.0)
16	-2026年版本计划：升级triton版本到triton3.4	16	+Version plan for 2026: Upgrade to Triton 3.4.
17	-- #### 里程碑	17	+- #### Milestones
18	-\| 里程碑 \| 重要特性更新情况 \| 状态 \|	18	+\| Milestone\| Important Update\| Status\|
19	\|------\|------\|------\|	19	\|------\|------\|------\|
20	-\| 2025.11.14 \| Triton-Ascend 3.2.0rc4预发布版本上线：<br>[扩展 tt.fp_to_fp 接口，新增对 FP8的类型转换支持](https://gitcode.com/Ascend/triton-ascend/pull/891) <br>[新增 scatter_ub_to_out 接口，支持从UB到GM的高效数据分散操作](https://gitcode.com/Ascend/triton-ascend/pull/864)\| ✅ \|	20	+\| 2025.11.14 \| The pre-release version Triton-Ascend 3.2.0rc4 is available.<br>[Extended the tt.fp_to_fp API to support conversion to the FP8 type.](https://gitcode.com/Ascend/triton-ascend/pull/891)<br>[Added the scatter_ub_to_out API to support efficient data scattering from the UB to the GM.](https://gitcode.com/Ascend/triton-ascend/pull/864)\| ✅ \|
21	-\| 2025.09.30 \| 完善Scan/Sort类Triton Python API，支持非连续访存，完成vLLM、sglang开源仓中重点Triton算子适配 \| ✅ \|	21	+\| 2025.09.30 \| Improved the Triton Python APIs of the Scan/Sort class, supported non-contiguous memory access, and completed the adaptation of key Triton operators in the vLLM and sglang open-source repositories.\| ✅ \|
22	-\| 2025.09.19 \| 支持Triton-Ascend [nightly包](https://test.pypi.org/project/triton-ascend/#history)提取 \| ✅ \|	22	+\| 2025.09.19 \| Supported the extraction of the Triton-Ascend [nightly package](https://test.pypi.org/project/triton-ascend/#history). \| ✅ \|
23	-\| 2025.08.15 \| 完善Atomic类Triton Python API支持，完成Flaggems开源仓重点Triton算子适配，提供Matmul等简单算子高性能实现参考用例 \| ✅ \|	23	+\| 2025.08.15 \| Improved the support for the Triton Python APIs of the Atomic class, completed the adaptation of key Triton operators in the Flaggems open-source repository, and provided reference cases for high-performance implementation of simple operators such as Matmul.\| ✅ \|
24	-\| 2025.06.30 \| 支持85% Triton Python API，支持连续访存，覆盖基本使用场景需求 \| ✅ \|	24	+\| 2025.06.30 \| Supported 85% of Triton Python APIs and contiguous memory access, covering basic application scenarios.\| ✅ \|
25	-\| 2025.05.20 \| Triton-Ascend开源，Gitcode代码仓Alive！ \| ✅ \|	25	+\| 2025.05.20 \| Triton-Ascend is open-source, and the GitCode code repository is alive!\| ✅ \|
26	-- #### 社区活动信息	26	+- #### Community Activities
27	-1. [会议日历](https://meeting.osinfra.cn/ascend)	27	+1. [Meeting calendar](https://meeting.osinfra.cn/ascend)
28	-2. [会议纪要看板]( https://etherpad-ascend.meeting.osinfra.cn/p/sig-AscendNPU-IR)	28	+2. [Meeting minutes dashboard](https://etherpad-ascend.meeting.osinfra.cn/p/sig-AscendNPU-IR)
29		29
30	-## 性能基准测试	30	+## Performance Benchmarking
31	-### 关键算子性能图表	31	+### Performance Charts of Key Operators
32	-选取经过性能优化后的关键算子FA、MM、Softmax作为示例。通过图表展示Triton算子与AscendC算子的性能差异，指标为加速比(`Speedup= AscendC_Duration_Time / Triton_Duration_Time`), [调优指南参考方法](./docs/zh/debug_guide/profiling.md)：	32	+The key operators FA, MM, and Softmax that have been optimized are selected as examples. The following charts show the performance differences between Triton operators and AscendC operators. The metric is the speedup ratio (`Speedup = AscendC_Duration_Time/Triton_Duration_Time`). For details, see the [Optimization Guide](./docs/en/debug_guide/profiling.md).
33		33
34	-- FA 性能图表：	34	+- FA performance chart:
35		35
36	-![FA 性能图表](docs/zh/figures/FA_Performance.png)	36	+![FA performance chart](docs/en/figures/FA_Performance.png)
37		37
38	-- MM 性能图表：	38	+- MM performance chart:
39		39
40	-![MM 性能图表](docs/zh/figures/MM_Performance.png)	40	+![MM performance chart](docs/en/figures/MM_Performance.png)
41		41
42	-- Softmax 性能图表：	42	+- Softmax performance chart:
43		43
44	-![Softmax 性能图表](docs/zh/figures/Softmax_Performance.png)	44	+![Softmax performance chart](docs/en/figures/Softmax_Performance.png)
45		45
46	-## 支持范围	46	+## Support
47		47
48		48
49	-- #### 硬件支持	49	+- #### Hardware Support
50	-Triton-Ascend 在昇腾 AI 产品支持使用，具体型号如下：	50	+Triton-Ascend is supported by Ascend AI products. The following table lists the product models.
51		51
52	-\| 产品系列 \| 产品型号 \|	52	+\| Product Series \| Product Model \|
53	\|----------------------------\|---------------------------------------\|	53	\|----------------------------\|---------------------------------------\|
54	-\| Atlas A3 训练系列产品 \| Atlas 800T A3 超节点服务器 \|	54	+\| Atlas A3 training products \| Atlas 800T A3 SuperNode server \|
55	-\| \| Atlas 900 A3 SuperPoD 超节点 \|	55	+\| \| Atlas 900 A3 SuperPoD server \|
56	-\| \| A200T A3 Box8 超节点服务器 \|	56	+\| \| A200T A3 Box8 SuperPoD server \|
57	-\| Atlas A3 推理系列产品 \| Atlas 800I A3 超节点服务器 \|	57	+\| Atlas A3 inference products \| Atlas 800I A3 SuperNode server \|
58	-\| Atlas A2 训练系列产品 \| Atlas 800T A2 训练服务器 \|	58	+\| Atlas A2 training products \| Atlas 800T A2 training server \|
59	-\| \| Atlas 900 A2 PoD 集群基础单元 \|	59	+\| \| Atlas 900 A2 PoD cluster basic unit \|
60	-\| \| Atlas 200T A2 Box16 异构子框 \|	60	+\| \| Atlas 200T A2 Box16 heterogeneous subrack \|
61	-\| Atlas A2 推理系列产品 \| Atlas 800I A2 推理服务器 \|	61	+\| Atlas A2 inference products \| Atlas 800I A2 inference server \|
62	-\| \| Atlas 300I A2 推理卡 \|	62	+\| \| Atlas 300I A2 inference card \|
63	-\| \| A200I A2 Box 异构组件 \|	63	+\| \| A200I A2 Box heterogeneous subrack \|
64		64
65	-- #### 兼容性	65	+- #### Compatibility
66		66
67	-支持操作系统：	67	+Supported OSs:
68	-Triton-Ascend 所支持的操作系统与 CANN 一致。请参考 CANN 官方文档，下载并安装适用于您操作系统的 CANN 版本。	68	+The OSs supported by Triton-Ascend are the same as those supported by CANN. Download and install the CANN version that is compatible with your OS. For details, see the official CANN documentation.
69		69
70	-CANN版本：	70	+CANN versions:
71		71
72	-- 商用版	72	+- Commercial versions
73		73
74	-\| Triton-Ascend版本 \| CANN商用版本 \| CANN发布日期 \|	74	+\| Triton-Ascend Version\| CANN Commercial Version\| Release Date\|
75	\|-------------------\|----------------------\|--------------------\|	75	\|-------------------\|----------------------\|--------------------\|
76	\| 3.2.0 \| CANN 8.5.0 \| 2026/01/16 \|	76	\| 3.2.0 \| CANN 8.5.0 \| 2026/01/16 \|
77	\| 3.2.0rc4 \| CANN 8.3.RC2 \| 2025/11/20 \|	77	\| 3.2.0rc4 \| CANN 8.3.RC2 \| 2025/11/20 \|
78	\| \| CANN 8.3.RC1 \| 2025/10/30 \|	78	\| \| CANN 8.3.RC1 \| 2025/10/30 \|
79		79
80	-- 社区版	80	+- Community versions
81		81
82	-\| Triton-Ascend版本 \| CANN社区版本 \| CANN发布日期 \|	82	+\| Triton-Ascend Version\| CANN Community Version\| Release Date\|
83	\|-------------------\|----------------------\|--------------------\|	83	\|-------------------\|----------------------\|--------------------\|
84	\| 3.2.0 \| CANN 8.5.0 \| 2026/01/16 \|	84	\| 3.2.0 \| CANN 8.5.0 \| 2026/01/16 \|
85	\| 3.2.0rc4 \| CANN 8.3.RC2 \| 2025/11/20 \|	85	\| 3.2.0rc4 \| CANN 8.3.RC2 \| 2025/11/20 \|
86	\| \| CANN 8.5.0.alpha001 \| 2025/11/12 \|	86	\| \| CANN 8.5.0.alpha001 \| 2025/11/12 \|
87	\| \| CANN 8.3.RC1 \| 2025/10/30 \|	87	\| \| CANN 8.3.RC1 \| 2025/10/30 \|
88		88
89	-## 入门指引	89	+## Getting Started
90		90
91	-- [快速开始](./docs/zh/quick_start.md)	91	+- [Quick Start](./docs/en/quick_start.md)
92		92
93	-- [架构设计与核心特性](./docs/zh/architecture_design_and_core_features.md)	93	+- [Architecture Design and Core Features](./docs/en/architecture_design_and_core_features.md)
94		94
95	-- [算子开发指南](./docs/zh/programming_guide.md)	95	+- [Operator Development Guide](./docs/en/programming_guide.md)
96		96
97	-- [算子迁移指南](./docs/zh/migration_guide/migrate_from_gpu.md)	97	+- [Operator Migration Guide](./docs/en/migration_guide/migrate_from_gpu.md)
98		98
99	-- [算子调试指南](./docs/zh/debug_guide/debugging.md#)	99	+- [Operator Debugging Guide](./docs/en/debug_guide/debugging.md#)
100		100
101	-- [性能调优指南](./docs/zh/debug_guide/profiling.md#)	101	+- [Performance Optimization Guide](./docs/en/debug_guide/profiling.md#)
102		102
103	-- [环境变量](docs/zh/environment_variable_reference.md)	103	+- [Environment Variables](docs/en/environment_variable_reference.md)
104		104
105	-## 常见问题	105	+## FAQ
106		106
107	-在使用Triton-Ascend时遇到的常见问题，详见 [FAQ](./docs/zh/FAQ.md#)	107	+For details about the FAQ encountered when using Triton-Ascend, see [FAQ](./docs/en/FAQ.md#).
108		108
109	-## 安全声明	109	+## Security Note
110		110
111	-我们重视开发者在使用Triton-Ascend时的信息安全，安全防护建议与相关信息请见 [安全声明](./SECURITYNOTE.md)	111	+We attach great importance to the information security of developers using Triton-Ascend. For details about the security protection suggestions and related information, see [Security Note](./SECURITYNOTE.md).
112		112
113	-## 许可证信息	113	+## License Information
114	-本项目代码与文档均采用 [MIT许可证](./LICENSE)	114	+The code and documents of this project are released under the [MIT License](./LICENSE).
115		115
116	-## 社区与贡献	116	+## Community and Contribution
117	-欢迎参与Triton-Ascend的开发及代码贡献，详情请参阅 [贡献指南](./CONTRIBUTING.zh.md)	117	+You are welcome to participate in the development and code contribution of Triton-Ascend. For details, see [Contribution Guide](./CONTRIBUTING.zh.md).

AREADME_zh.md+117-0

@@ -0,0 +1,117 @@
	1	+# Triton-Ascend
	2	+
	3	+## 项目简介与价值主张
	4	+Triton-Ascend是面向昇腾平台构建的Triton编译框架，旨在让Triton代码能够在昇腾硬件上高效运行。
	5	+- #### 核心价值说明
	6	+Triton是近几年来受到开发者青睐的Python化编译框架。开发者仅需关注Tile/Block的切分方式以及基于Tile/Block的运算逻辑，编译器将在Triton代码的编译过程中结合底层硬件特点自动完成内存分配、数据搬运、数据计算、流水并行等，因此，算子的开发难度大幅降低、开发效率显著提升。
	7	+Triton-Ascend将Triton编译栈适配到华为昇腾NPU上，在Triton的基础上提供一系列针对性的优化，使Triton代码能够编译后在昇腾硬件上高效运行。
	8	+目前，Triton-Ascend仍在持续完善中，我们将不断提升Triton Python API完备度、数据类型支持度、访存方式灵活性等，并持续优化编译器的自动优化能力，提升Triton-Ascend整体的功能与性能泛化性。
	9	+- #### 昇腾生态定位
	10	+Triton-Ascend编译框架打通了Triton与昇腾硬件之间的壁垒，使熟悉Triton框架的开发者可以更有效率地使用昇腾NPU。它通过提供通用、高效的算子开发范式，为昇腾软件栈补齐了敏捷开发的关键一环，极大丰富了昇腾的算子库和上层应用生态。
	11	+
	12	+## 最新动态与里程碑
	13	+- #### 近期版本更新
	14	+当前版本：[Triton-Ascend 3.2.0](https://pypi.org/project/triton-ascend/)
	15	+配套CANN版本：[昇腾CANN社区版8.5.0](https://www.hiascend.com/developer/download/community/result?module=cann&cann=8.5.0)
	16	+2026年版本计划：升级triton版本到triton3.4
	17	+- #### 里程碑
	18	+\| 里程碑 \| 重要特性更新情况 \| 状态 \|
	19	+\|------\|------\|------\|
	20	+\| 2025.11.14 \| Triton-Ascend 3.2.0rc4预发布版本上线：<br>[扩展 tt.fp_to_fp 接口，新增对 FP8的类型转换支持](https://gitcode.com/Ascend/triton-ascend/pull/891) <br>[新增 scatter_ub_to_out 接口，支持从UB到GM的高效数据分散操作](https://gitcode.com/Ascend/triton-ascend/pull/864)\| ✅ \|
	21	+\| 2025.09.30 \| 完善Scan/Sort类Triton Python API，支持非连续访存，完成vLLM、sglang开源仓中重点Triton算子适配 \| ✅ \|
	22	+\| 2025.09.19 \| 支持Triton-Ascend [nightly包](https://test.pypi.org/project/triton-ascend/#history)提取 \| ✅ \|
	23	+\| 2025.08.15 \| 完善Atomic类Triton Python API支持，完成Flaggems开源仓重点Triton算子适配，提供Matmul等简单算子高性能实现参考用例 \| ✅ \|
	24	+\| 2025.06.30 \| 支持85% Triton Python API，支持连续访存，覆盖基本使用场景需求 \| ✅ \|
	25	+\| 2025.05.20 \| Triton-Ascend开源，Gitcode代码仓Alive！ \| ✅ \|
	26	+- #### 社区活动信息
	27	+1. [会议日历](https://meeting.osinfra.cn/ascend)
	28	+2. [会议纪要看板]( https://etherpad-ascend.meeting.osinfra.cn/p/sig-AscendNPU-IR)
	29	+
	30	+## 性能基准测试
	31	+### 关键算子性能图表
	32	+选取经过性能优化后的关键算子FA、MM、Softmax作为示例。通过图表展示Triton算子与AscendC算子的性能差异，指标为加速比(`Speedup= AscendC_Duration_Time / Triton_Duration_Time`), [调优指南参考方法](./docs/zh/debug_guide/profiling.md)：
	33	+
	34	+- FA 性能图表：
	35	+
	36	+![FA 性能图表](docs/zh/figures/FA_Performance.png)
	37	+
	38	+- MM 性能图表：
	39	+
	40	+![MM 性能图表](docs/zh/figures/MM_Performance.png)
	41	+
	42	+- Softmax 性能图表：
	43	+
	44	+![Softmax 性能图表](docs/zh/figures/Softmax_Performance.png)
	45	+
	46	+## 支持范围
	47	+
	48	+
	49	+- #### 硬件支持
	50	+Triton-Ascend 在昇腾 AI 产品支持使用，具体型号如下：
	51	+
	52	+\| 产品系列 \| 产品型号 \|
	53	+\|----------------------------\|---------------------------------------\|
	54	+\| Atlas A3 训练系列产品 \| Atlas 800T A3 超节点服务器 \|
	55	+\| \| Atlas 900 A3 SuperPoD 超节点 \|
	56	+\| \| A200T A3 Box8 超节点服务器 \|
	57	+\| Atlas A3 推理系列产品 \| Atlas 800I A3 超节点服务器 \|
	58	+\| Atlas A2 训练系列产品 \| Atlas 800T A2 训练服务器 \|
	59	+\| \| Atlas 900 A2 PoD 集群基础单元 \|
	60	+\| \| Atlas 200T A2 Box16 异构子框 \|
	61	+\| Atlas A2 推理系列产品 \| Atlas 800I A2 推理服务器 \|
	62	+\| \| Atlas 300I A2 推理卡 \|
	63	+\| \| A200I A2 Box 异构组件 \|
	64	+
	65	+- #### 兼容性
	66	+
	67	+支持操作系统：
	68	+Triton-Ascend 所支持的操作系统与 CANN 一致。请参考 CANN 官方文档，下载并安装适用于您操作系统的 CANN 版本。
	69	+
	70	+CANN版本：
	71	+
	72	+- 商用版
	73	+
	74	+\| Triton-Ascend版本 \| CANN商用版本 \| CANN发布日期 \|
	75	+\|-------------------\|----------------------\|--------------------\|
	76	+\| 3.2.0 \| CANN 8.5.0 \| 2026/01/16 \|
	77	+\| 3.2.0rc4 \| CANN 8.3.RC2 \| 2025/11/20 \|
	78	+\| \| CANN 8.3.RC1 \| 2025/10/30 \|
	79	+
	80	+- 社区版
	81	+
	82	+\| Triton-Ascend版本 \| CANN社区版本 \| CANN发布日期 \|
	83	+\|-------------------\|----------------------\|--------------------\|
	84	+\| 3.2.0 \| CANN 8.5.0 \| 2026/01/16 \|
	85	+\| 3.2.0rc4 \| CANN 8.3.RC2 \| 2025/11/20 \|
	86	+\| \| CANN 8.5.0.alpha001 \| 2025/11/12 \|
	87	+\| \| CANN 8.3.RC1 \| 2025/10/30 \|
	88	+
	89	+## 入门指引
	90	+
	91	+- [快速开始](./docs/zh/quick_start.md)
	92	+
	93	+- [架构设计与核心特性](./docs/zh/architecture_design_and_core_features.md)
	94	+
	95	+- [算子开发指南](./docs/zh/programming_guide.md)
	96	+
	97	+- [算子迁移指南](./docs/zh/migration_guide/migrate_from_gpu.md)
	98	+
	99	+- [算子调试指南](./docs/zh/debug_guide/debugging.md#)
	100	+
	101	+- [性能调优指南](./docs/zh/debug_guide/profiling.md#)
	102	+
	103	+- [环境变量](docs/zh/environment_variable_reference.md)
	104	+
	105	+## 常见问题
	106	+
	107	+在使用Triton-Ascend时遇到的常见问题，详见 [FAQ](./docs/zh/FAQ.md#)
	108	+
	109	+## 安全声明
	110	+
	111	+我们重视开发者在使用Triton-Ascend时的信息安全，安全防护建议与相关信息请见 [安全声明](./SECURITYNOTE.md)
	112	+
	113	+## 许可证信息
	114	+本项目代码与文档均采用 [MIT许可证](./LICENSE)
	115	+
	116	+## 社区与贡献
	117	+欢迎参与Triton-Ascend的开发及代码贡献，详情请参阅 [贡献指南](./CONTRIBUTING.zh.md)

MSECURITYNOTE.md+43-43

	@@ -1,63 +1,63 @@
1	-# Triton-Ascend 安全声明	1	+# Triton-Ascend Security Note
2		2
3	-## 系统安全加固	3	+## System Security Hardening
4		4
5	-建议用户在系统中配置开启ASLR（级别2 ），又称全随机地址空间布局随机化，可参考以下方式进行配置：	5	+You are advised to enable the address space layout randomization (ASLR) (level 2) in the system. You can perform the following operation to enable it:
6		6
7	echo 2 > /proc/sys/kernel/randomize_va_space	7	echo 2 > /proc/sys/kernel/randomize_va_space
8		8
9	-## 运行用户建议	9	+## Suggestions on Running Users
10		10
11	-出于安全性及权限最小化角度考虑，不建议通过root等管理员类型账户使用Triton-Ascend。	11	+To ensure security and minimize permissions, you are advised not to use administrator accounts such as root.
12		12
13	-## 文件权限控制	13	+## File Permission Control
14		14
15	-1. 建议用户对个人的隐私数据、商业资产等敏感文件做好权限控制等安全措施，设定的权限建议参考[文件权限参考](#文件权限参考)进行设置。	15	+1. You are advised to take security measures such as permission control on sensitive files, such as personal privacy data and business assets. For details about how to set the permissions, see the "File Permission Reference" section.
16		16
17	-2. 用户安装和使用过程需要做好权限控制，建议参考[文件权限参考](#文件权限参考)进行设置。	17	+2. During the installation and use, you are advised to control the permissions. For details about how to set the permissions, see [File Permission Reference](#file-permission-reference).
18		18
19		19
20	-##### 文件权限参考	20	+##### File permission reference
21		21
22	-\| 类型 \| Linux权限参考最大值 \|	22	+\| Type \| Maximum Permission in Linux \|
23	\|----------------------------------- \|-----------------------\|	23	\|----------------------------------- \|-----------------------\|
24	-\| 用户主目录 \| 750（rwxr-x---） \|	24	+\| Home directory \| 750 (rwxr-x---) \|
25	-\| 程序文件（含脚本文件、库文件等） \| 550（r-xr-x---） \|	25	+\| Program files (including scripts and libraries) \| 550 (r-xr-x---) \|
26	-\| 程序文件目录 \| 550（r-xr-x---） \|	26	+\| Program file directory \| 550 (r-xr-x---) \|
27	-\| 配置文件 \| 640（rw-r-----） \|	27	+\| Configuration files \| 640 (rw-r-----) \|
28	-\| 配置文件目录 \| 750（rwxr-x---） \|	28	+\| Configuration file directory \| 750 (rwxr-x---) \|
29	-\| 日志文件（记录完毕或者已经归档） \| 440（r--r-----） \|	29	+\| Log files (recorded or archived) \| 440 (r--r-----) \|
30	-\| 日志文件（正在记录） \| 640（rw-r-----） \|	30	+\| Log files (being recorded) \| 640 (rw-r-----) \|
31	-\| 日志文件目录 \| 750（rwxr-x---） \|	31	+\| Log file directory \| 750 (rwxr-x---) \|
32	-\| Debug文件 \| 640（rw-r-----） \|	32	+\| Debug files \| 640 (rw-r-----) \|
33	-\| Debug文件目录 \| 750（rwxr-x---） \|	33	+\| Debug file directory \| 750 (rwxr-x---) \|
34	-\| 临时文件目录 \| 750（rwxr-x---） \|	34	+\| Temporary file directory \| 750 (rwxr-x---) \|
35	-\| 维护升级文件目录 \| 770（rwxrwx---） \|	35	+\| Maintenance and upgrade file directory \| 770 (rwxrwx---) \|
36	-\| 业务数据文件 \| 640（rw-r-----） \|	36	+\| Service data files \| 640 (rw-r-----) \|
37	-\| 业务数据文件目录 \| 750（rwxr-x---） \|	37	+\| Service data file directory \| 750 (rwxr-x---) \|
38	-\| 密钥组件、私钥、证书、密文文件目录 \| 700（rwx------） \|	38	+\| Key component, private key, certificate, and ciphertext file directory \| 700 (rwx------) \|
39	-\| 密钥组件、私钥、证书、加密密文 \| 600（rw-------） \|	39	+\| Key components, private keys, certificates, and ciphertext files \| 600 (rw-------) \|
40	-\| 加解密接口、加解密脚本 \| 500（r-x------） \|	40	+\| APIs and script files for encryption and decryption \| 500 (r-x------) \|
41		41
42		42
43	-## 构建安全声明	43	+## Build Security Statement
44		44
45	-Triton-Ascend支持源码编译安装，在编译时会下载依赖第三方库并执行构建shell脚本，在编译过程中会产生临时程序文件和编译目录。用户可根据需要自行对源代码目录内的文件进行权限管控降低安全风险。	45	+Triton-Ascend can be installed through source code compilation. During the compilation process, dependent third-party libraries are downloaded and the shell build script is executed. This results in the generation of temporary program files and compilation directories. You can control permissions on files in the source code directory as required to prevent security risks.
46		46
47	-## 公网地址声明	47	+## Public IP Address Statement
48		48
49	-在Triton-Ascend的配置文件和脚本中存在[公网地址](#公网地址)	49	+Public IP addresses are used in the configuration files and scripts of Triton-Ascend. For details, see the "Public IP Addresses" section.
50		50
51	-##### 公网地址	51	+##### Public IP addresses
52	-\| 类型 \| 开源代码地址 \| 文件名 \| 公网IP地址/公网URL地址/域名/邮箱地址 \| 用途说明 \|	52	+\| Type \| Open-Source Code Address \| File Name \| Public IP Address/Public URL/Domain Name/Email Address \| Description \|
53	\|----------\|------------------------------------------------------------------------------------------------\|-------------------------------------------\|------------------------------------------------------------------------------------------------------\|-----------------------------------\|	53	\|----------\|------------------------------------------------------------------------------------------------\|-------------------------------------------\|------------------------------------------------------------------------------------------------------\|-----------------------------------\|
54	-\| 开源引入 \| https://github.com/triton-lang/triton.git \| .gitmodules \| https://github.com/triton-lang/triton.git \| Triton源码仓地址 \|	54	+\| Introduced by open source\| https://github.com/triton-lang/triton.git \| .gitmodules \| https://github.com/triton-lang/triton.git \| Address of the Triton source code repository\|
55	-\| 开源引入 \| https://gitcode.com/Ascend/AscendNPU-IR.git \| .gitmodules \| https://gitcode.com/Ascend/AscendNPU-IR.git \| AscendNPU IR源码仓地址 \|	55	+\| Introduced by open source\| https://gitcode.com/Ascend/AscendNPU-IR.git \| .gitmodules \| https://gitcode.com/Ascend/AscendNPU-IR.git \| AscendNPU IR source code repository address\|
56	-\| 自研 \| 不涉及 \| docker/devdocker/setup_triton-ascend_dev.sh \| https://gitcode.com/Ascend/triton-ascend.git \| Triton-Ascend源码仓地址 \|	56	+\| Self-developed \| N/A \| docker/devdocker/setup_triton-ascend_dev.sh \| https://gitcode.com/Ascend/triton-ascend.git \| Address of the Triton-Ascend source code repository \|
57	-\| 自研 \| 不涉及 \| ascend/examples/generalization_cases/run_daily.sh & scripts/prepare_build.sh \| https://gitee.com/shijingchang/triton.git \| 构建依赖代码仓 \|	57	+\| Self-developed \| N/A \| ascend/examples/generalization_cases/run_daily.sh & scripts/prepare_build.sh \| https://gitee.com/shijingchang/triton.git \| Build dependency code repository \|
58	-\| 自研 \| 不涉及 \| setup.py \| https://gitcode.com/Ascend/triton-ascend/ \| Triton-Ascend源码仓地址 \|	58	+\| Self-developed \| N/A \| setup.py \| https://gitcode.com/Ascend/triton-ascend/ \| Address of the Triton-Ascend source code repository\|
59	-\| 开源引入 \| https://gitclone.com \| scripts/prepare_build.sh \| https://gitclone.com/github.com/llvm/llvm-project.git \| 依赖的llvm源码仓 \|	59	+\| Introduced by open source\| https://gitclone.com \| scripts/prepare_build.sh \| https://gitclone.com/github.com/llvm/llvm-project.git \| LLVM source code repository \|
60	-\| 开源引入 \| https://repo.huaweicloud.com \| scripts/prepare_build.sh \| https://repo.huaweicloud.com/repository/pypi/simple \| 用于配置pybind11下载链接 \|	60	+\| Introduced by open source\| https://repo.huaweicloud.com \| scripts/prepare_build.sh \| https://repo.huaweicloud.com/repository/pypi/simple \| Used to configure the pybind11 download link.\|
61	-\| 开源引入 \| https://pypi.tuna.tsinghua.edu.cn \| docker/devdocker/triton-ascend_dev.dockerfile \| https://pypi.tuna.tsinghua.edu.cn/simple \| python pip源配置 \|	61	+\| Introduced by open source\| https://pypi.tuna.tsinghua.edu.cn \| docker/devdocker/triton-ascend_dev.dockerfile \| https://pypi.tuna.tsinghua.edu.cn/simple \| Python pip source configuration \|
62	-\| 开源引入 \| https://triton-ascend-artifacts.obs.myhuaweicloud.com \| setup.py \|https://triton-ascend-artifacts.obs.myhuaweicloud.com/llvm-builds/{name}.tar.gz \| 用于下载预编译的LLVM工具 \|	62	+\| Introduced by open source\| https://triton-ascend-artifacts.obs.myhuaweicloud.com \| setup.py \|https://triton-ascend-artifacts.obs.myhuaweicloud.com/llvm-builds/{name}.tar.gz \| Used to download the prepared LLVM tool.\|
63	-\| 开源引入 \| https://bootstrap.pypa.io/get-pip.py \| docker/develop_env.dockerfile \|https://bootstrap.pypa.io/get-pip.py \| 用于自动化安装pip \|	63	+\| Introduced by open source\| https://bootstrap.pypa.io/get-pip.py \| docker/develop_env.dockerfile \|https://bootstrap.pypa.io/get-pip.py \| Used to automatically install pip.\|

ASECURITYNOTE_zh.md+63-0

@@ -0,0 +1,63 @@
	1	+# Triton-Ascend 安全声明
	2	+
	3	+## 系统安全加固
	4	+
	5	+建议用户在系统中配置开启ASLR（级别2 ），又称全随机地址空间布局随机化，可参考以下方式进行配置：
	6	+
	7	+ echo 2 > /proc/sys/kernel/randomize_va_space
	8	+
	9	+## 运行用户建议
	10	+
	11	+出于安全性及权限最小化角度考虑，不建议通过root等管理员类型账户使用Triton-Ascend。
	12	+
	13	+## 文件权限控制
	14	+
	15	+1. 建议用户对个人的隐私数据、商业资产等敏感文件做好权限控制等安全措施，设定的权限建议参考[文件权限参考](#文件权限参考)进行设置。
	16	+
	17	+2. 用户安装和使用过程需要做好权限控制，建议参考[文件权限参考](#文件权限参考)进行设置。
	18	+
	19	+
	20	+##### 文件权限参考
	21	+
	22	+\| 类型 \| Linux权限参考最大值 \|
	23	+\|----------------------------------- \|-----------------------\|
	24	+\| 用户主目录 \| 750（rwxr-x---） \|
	25	+\| 程序文件（含脚本文件、库文件等） \| 550（r-xr-x---） \|
	26	+\| 程序文件目录 \| 550（r-xr-x---） \|
	27	+\| 配置文件 \| 640（rw-r-----） \|
	28	+\| 配置文件目录 \| 750（rwxr-x---） \|
	29	+\| 日志文件（记录完毕或者已经归档） \| 440（r--r-----） \|
	30	+\| 日志文件（正在记录） \| 640（rw-r-----） \|
	31	+\| 日志文件目录 \| 750（rwxr-x---） \|
	32	+\| Debug文件 \| 640（rw-r-----） \|
	33	+\| Debug文件目录 \| 750（rwxr-x---） \|
	34	+\| 临时文件目录 \| 750（rwxr-x---） \|
	35	+\| 维护升级文件目录 \| 770（rwxrwx---） \|
	36	+\| 业务数据文件 \| 640（rw-r-----） \|
	37	+\| 业务数据文件目录 \| 750（rwxr-x---） \|
	38	+\| 密钥组件、私钥、证书、密文文件目录 \| 700（rwx------） \|
	39	+\| 密钥组件、私钥、证书、加密密文 \| 600（rw-------） \|
	40	+\| 加解密接口、加解密脚本 \| 500（r-x------） \|
	41	+
	42	+
	43	+## 构建安全声明
	44	+
	45	+Triton-Ascend支持源码编译安装，在编译时会下载依赖第三方库并执行构建shell脚本，在编译过程中会产生临时程序文件和编译目录。用户可根据需要自行对源代码目录内的文件进行权限管控降低安全风险。
	46	+
	47	+## 公网地址声明
	48	+
	49	+在Triton-Ascend的配置文件和脚本中存在[公网地址](#公网地址)
	50	+
	51	+##### 公网地址
	52	+\| 类型 \| 开源代码地址 \| 文件名 \| 公网IP地址/公网URL地址/域名/邮箱地址 \| 用途说明 \|
	53	+\|----------\|------------------------------------------------------------------------------------------------\|-------------------------------------------\|------------------------------------------------------------------------------------------------------\|-----------------------------------\|
	54	+\| 开源引入 \| https://github.com/triton-lang/triton.git \| .gitmodules \| https://github.com/triton-lang/triton.git \| Triton源码仓地址 \|
	55	+\| 开源引入 \| https://gitcode.com/Ascend/AscendNPU-IR.git \| .gitmodules \| https://gitcode.com/Ascend/AscendNPU-IR.git \| AscendNPU IR源码仓地址 \|
	56	+\| 自研 \| 不涉及 \| docker/devdocker/setup_triton-ascend_dev.sh \| https://gitcode.com/Ascend/triton-ascend.git \| Triton-Ascend源码仓地址 \|
	57	+\| 自研 \| 不涉及 \| ascend/examples/generalization_cases/run_daily.sh & scripts/prepare_build.sh \| https://gitee.com/shijingchang/triton.git \| 构建依赖代码仓 \|
	58	+\| 自研 \| 不涉及 \| setup.py \| https://gitcode.com/Ascend/triton-ascend/ \| Triton-Ascend源码仓地址 \|
	59	+\| 开源引入 \| https://gitclone.com \| scripts/prepare_build.sh \| https://gitclone.com/github.com/llvm/llvm-project.git \| 依赖的llvm源码仓 \|
	60	+\| 开源引入 \| https://repo.huaweicloud.com \| scripts/prepare_build.sh \| https://repo.huaweicloud.com/repository/pypi/simple \| 用于配置pybind11下载链接 \|
	61	+\| 开源引入 \| https://pypi.tuna.tsinghua.edu.cn \| docker/devdocker/triton-ascend_dev.dockerfile \| https://pypi.tuna.tsinghua.edu.cn/simple \| python pip源配置 \|
	62	+\| 开源引入 \| https://triton-ascend-artifacts.obs.myhuaweicloud.com \| setup.py \|https://triton-ascend-artifacts.obs.myhuaweicloud.com/llvm-builds/{name}.tar.gz \| 用于下载预编译的LLVM工具 \|
	63	+\| 开源引入 \| https://bootstrap.pypa.io/get-pip.py \| docker/develop_env.dockerfile \|https://bootstrap.pypa.io/get-pip.py \| 用于自动化安装pip \|

Adocs/en/FAQ.md+52-0

@@ -0,0 +1,52 @@
	1	+# Triton-Ascend FAQ
	2	+
	3	+## 1. Installation and Environment Configuration
	4	+
	5	+Q: How can I correctly install Triton-Ascend? Is it possible to install it directly using pip?
	6	+
	7	+A: You can directly use pip to install it.
	8	+```Python
	9	+pip install triton-ascend
	10	+```
	11	+Q: Can Triton-Ascend be used on non-Ascend hardware (such as CUDA AMD)?
	12	+
	13	+A: No. Triton-Ascend can be used only in the Ascend NPU hardware environment.
	14	+
	15	+## 2. Accuracy and Numerical Consistency Issues
	16	+
	17	+Q: How can I troubleshoot the inconsistency between the NPU running result and the PyTorch/CPU/GPU reference result?
	18	+
	19	+A: For details, see [07_accuracy_comparison_example.md](../en/examples/07_accuracy_comparison_example.md).
	20	+For details about the debugging method, see [Debugging in Interpreter Mode](./debug_guide/debugging.md#4-interpreter-mode).
	21	+
	22	+## 3. Error Code and Exception Handling
	23	+
	24	+Q: Why is the error message "MLIRCompilationError" displayed during kernel compilation? How can I locate the failed pass?
	25	+
	26	+A: For details, see [Compilation Error Debugging](./debug_guide/debugging.md#52-compilation-errors-debugging).
	27	+
	28	+## 4. Debugging and Logging
	29	+
	30	+Q: How can I enable detailed log output? Where is the output of TRITON_DEBUG=1?
	31	+
	32	+A: You can use TRITON_DEBUG=1 to obtain detailed dump files for debugging. For details, see [Dump Files](./debug_guide/debugging.md#32-dump-files).
	33	+
	34	+Q: Can I print the intermediate tensor value in the kernel? Is tl.device_print available?
	35	+
	36	+A: You can use tl.device_print to print the tensor in the kernel. For details, see [Debugging by Printing](./debug_guide/debugging.md#51-debugging-by-printing).
	37	+
	38	+## 5. Development and Contributions
	39	+
	40	+Q: How can I build and test Triton-Ascend locally?
	41	+
	42	+A: For details about the local build and test methods, see [Installing Triton-Ascend Using the Source Code](./installation_guide.md#installing-triton-ascend-using-the-source-code).
	43	+
	44	+Q: What CI checks are required for submitting a PR?
	45	+
	46	+A: The CI checks for a PR include: coding security and specifications check, open-source code check, malicious code check, compilation and building, and developer testing.
	47	+
	48	+## 6. Performance Optimization
	49	+
	50	+Q: Is there any performance analysis tool (profiler) available?
	51	+
	52	+A: There is an integrated performance analysis tool (profiler). For details, see [Operator Performance Optimization Methods](./debug_guide/profiling.md).

Adocs/en/architecture_design_and_core_features.md+221-0

@@ -0,0 +1,221 @@
	1	+# Architecture Design and Core Features
	2	+
	3	+## 1. Logical Architecture
	4	+![Image](./figures/architectural_diagram.png)
	5	+Triton-Ascend Architecture Description
	6	+
	7	+Core components:
	8	+- `Ascend language extension`: Triton language extension for Ascend
	9	+- `compiler`: Triton compiler for Ascend
	10	+- `driver`: Ascend device driver API
	11	+
	12	+Component functions:
	13	+
	14	+- `Ascend language extension`
	15	+ Syntax and semantic extensions for the Ascend NPU architecture are introduced based on the standard Triton language.
	16	+
	17	+- `compiler`
	18	+ Receives the Triton IR `(TTIR)` file generated by the upper-layer Triton compiler and performs a series of transformations to adapt to Ascend hardware.
	19	+ ```
	20	+ Triton IR → Linalg IR → AscendNPU IR → triton_xxx_kernel.o
	21	+ ```
	22	+ Converts the Triton IR into the Linalg IR, and then generates the executable binary file `triton_xxx_kernel.o` for the Ascend NPU through the BiSheng Compiler.
	23	+
	24	+- `driver`
	25	+ Provides the interconnection capability between the Triton runtime and the Ascend software stack (CANN), and loads the executable kernel file `triton_xxx_kernel.o` generated by the BiSheng Compiler on the device side.
	26	+
	27	+
	28	+## 2. Code Structure
	29	+### 2.1 Code Structure Principles
	30	+
	31	+This project extends the support for Huawei Ascend NPU (using the CANN software stack) based on the standard Triton. The overall design complies with the following code principles:
	32	+
	33	+> - If the modification is target independent, it should be retained in the Triton core part (such as general modifications to the language and runtime).
	34	+> - If the modification is target affinitive, it should be placed in the Triton-Ascend part.
	35	+
	36	+
	37	+### 2.2 Directory Structure and Function Description
	38	+
	39	+`include/` and `lib/`
	40	+- Content: MLIR Passes, dialects, and related tools for Ascend NPUs.
	41	+- Description: Represents and optimizes Ascend-specific computational graphs in the MLIR compilation process.
	42	+
	43	+`libdevice.py`
	44	+- Content: `libdevice` API adaptable to Ascend NPUs.
	45	+- Description: Provides underlying implementation support for the Ascend NPU hardware, which is called by Triton operators.
	46	+
	47	+
	48	+`backend/compiler.py`
	49	+- Content: Main entry of the `triton-ascend` compiler.
	50	+- Description: Compiles the high-level DSL code of Triton into an executable binary file (such as the`.o` file) that can be executed on the Ascend NPU.
	51	+
	52	+`backend/driver.py`
	53	+- Content: `triton-ascend` driver module.
	54	+- Description: Loads and starts the compiled executable binary file.
	55	+
	56	+
	57	+## 3. Modules
	58	+
	59	+### 3.1 Triton Core Enhancement
	60	+
	61	+#### 3.1.1 Language Extension
	62	+
	63	+\| No.\| Operator \| Description \|
	64	+\| :--- \| :--------------------------------------------- \| :--------------------------------- \|
	65	+\| 1 \| `tl.insert_slice(full, src, offsets, sizes, strides)` \| Inserts a tensor into another tensor according to the specified offset, size, and stride.<br>Returns: target tensor.<br>full: target tensor. The source tensor will be inserted into this tensor.<br>src: source tensor.<br>offsets: offset (integer tuple) on the target tensor.<br>sizes: size (integer tuple) on the source tensor.<br>strides: stride (integer tuple) on the target tensor.\|
	66	+\| 2 \| `tl.extract_slice(full, offsets, sizes, strides)` \| Extracts a slice tensor from another tensor according to the specified offset, size, and stride.<br>Returns: slice tensor.<br>full: source tensor. The slice is extracted from this tensor.<br>offsets: offset (integer tuple) on the source tensor.<br>sizes: size (integer tuple) of the slice tensor.<br>strides: stride (integer tuple) on the source tensor. \|
	67	+\| 3 \| `tl.get_element(source, offset)` \| Reads a tensor with dimensions and returns a single element at the specified offset.<br>source: source tensor.<br>offset: offset (integer tuple) of the element to be extracted. \|
	68	+
	69	+
	70	+## 3.2 Triton-Ascend
	71	+
	72	+### 3.2.1 Compiler Options
	73	+
	74	+\|No.\| NPU Option \| Hardware Platform \| Description\|
	75	+\| --- \| --------------------------------------------- \| ---------- \| ----- \|
	76	+\| 1 \| multibuffer \| NPU \| Autotune option. It enables or disables the ping-pong pipeline.\|
	77	+\| 2 \| enable_auto_bind_sub_block \| NPU \| Autotune option (CV-fused kernels only). It enables or disables auto-binding of sub-blocks.\|
	78	+\| 3 \| enable_hivm_auto_cv_balance \| NPU \| Autotune option (CV-fused kernels only). It enables or disables automatic CV balancing.\|
	79	+\| 4 \| sync_solver \| NPU \| Autotune option (CV-fused kernels only). It enables or disables the synchronization solver. \|
	80	+\| 5 \| unit_flag \| NPU \| Autotune option. It enables or disables the sync unit flag.\|
	81	+\| 6 \| inject_barrier_all \| NPU \| Autotune option. It enables or disables automatic injection of barriers for all operations.\|
	82	+\| 7 \| inject_block_all \| NPU \| Autotune option. It enables or disables automatic injection of blocks for all operations.\|
	83	+\| 8 \| limit_auto_multi_buffer_only_for_local_buffer \| NPU \| Autotune option. It restricts automatic multi-buffering only to local buffers.\|
	84	+\| 9 \| limit_auto_multi_buffer_of_local_buffer \| NPU \| Autotune option. It enables or disables automatic multi-buffering for local buffers.\|
	85	+\| 10 \| set_workspace_multibuffer \| NPU \| Autotune option. It enables or disables multi-buffering for the workspace.\|
	86	+\| 11 \| tile_mix_vector_loop \| NPU \| Autotune option (CV-fused kernels only). It enables or disables tiling for vector loops.\|
	87	+\| 12 \| tile_mix_cube_loop \| NPU \| Autotune option (CV-fused kernels only). It enables or disables tiling for cube loops.\|
	88	+\| 13 \| disable_auto_inject_block_sync \| NPU \| Autotune option (CV-fused kernels only). It enables or disables automatic injection of block synchronizations.\|
	89	+\| 14 \| stream \| NPU \| (Optional) Informs the compiler about the NPU stream to use.\|
	90	+\| 15 \| enable_linearize \| NPU \| Autotune option. It enables or disables the linearization pass.\|
	91	+\| 16 \| enable_nd2nz_on_vector \| NPU \| Autotune option (CV-fused kernels only). It enables or disables the ND (n-dimensional) to NZ (non-zero) layout transformation.\|
	92	+
	93	+### 3.2.2 SIMD Compiler
	94	+
	95	+
	96	+\| No.\| Pass \| Purpose \| IR Conversion \|
	97	+\| ------ \| ---------------------- \|----------------------------------------------------------------------\| ----------------------- \|
	98	+\| 1 \| triton-to-structured \| linearize \| ttir->ttir \|
	99	+\| 2 \| triton-to-unstructured \| convert indirect axis to loop \| ttir->ttir \|
	100	+\| 3 \| triton-to-linalg \| memory/reduction/view/creation/math/arith/linear algebra to linalgir \| ttir->linalgir \|
	101	+\| 4 \| triton-to-other \| ttir->hivm/hfusion/llvm \| ttir->hivm/hfusion/llvm \|
	102	+
	103	+#### 3.2.2.1 TritonToStructured
	104	+
	105	+The integer division and modulo operations in pointer expressions and mask expressions are converted to tensor operations to regenerate OPs such as load and store.
	106	+
	107	+\| Converter \| Description \| Limitation\|
	108	+\| ------------------------ \| -------------------------- \| ------------------------- \|
	109	+\| RewriteAddPtrOp \| Analyzes the pointer expression (`AddPtrOp`) in operations such as `tl.load` and `tl.store`. The original pointer offset is calculated and modeled into a `PtrState` object that contains the specific offset information of each dimension (axis). For example, for an expression like `ptr + x // 1024 * 4096 + x % 1024 * 4 + y`, the contributions and relationships of the `x` and `y` axes are analyzed. \| 1. The original iteration axis (such as `x`) must be exactly divisible by the split axis (such as `1024`).<br>2. The size of the external `XBLOCK` must be an integer multiple or divisor of the split axis `divisor`. \|
	110	+\| CreateAddpr \| Reconstructs a new `AddPtrOp` pointer calculation operation based on the analyzed `PtrState` object to eliminate the integer division (`//`) and modulo (`%`) operations in the original expression. \| It depends on the valid `PtrState` object successfully generated by `RewriteAddPtrOp`. \|
	111	+\| RewriteLoadOp \| Analyzes the `mask` expression in the `tl.load` operation and decomposes and models the complex mask conditions involving integer division and modulo into a `MaskState` object that contains the boundary information of each dimension. For example, for `mask = x // 1024 < 8 and x % 1024 < 1024 and y < 4`, an independent constraint condition of each dimension is obtained through analysis. \| 1. The original iteration axis (such as `x`) must be exactly divisible by the split axis (such as `1024`).<br>2. The size of the external `XBLOCK` must be an integer multiple or divisor of the split axis `divisor`. \|
	112	+\| BuildMask \| Reconstructs a new `mask` expression based on the analyzed `MaskState` object to eliminate the integer division (`//`) and modulo (`%`) operations in the original expression. \| Only the `MaskState` object generated by `RewriteLoadOp` or `RewriteStoreOp` is processed. Any complex and non-standardized mask expressions cannot be processed. \|
	113	+\| CreateLoad \| Re-creates/Replaces the original `tl.load` operation using the new pointer expression generated by `CreateAddpr` and the new mask expression generated by `BuildMask` to complete instruction rewriting. \| The prerequisite steps, such as `RewriteAddPtrOp`, `CreateAddpr`, `RewriteLoadOp`, `BuildMask`, must have been successfully executed. \|
	114	+\| RewriteStoreOp \| Analyzes the `mask` expression in the `tl.store` operation. Its function is similar to that of `RewriteLoadOp`. It decomposes and models the complex mask conditions involving integer division and modulo into a `MaskState` object. \| Same as `RewriteLoadOp`. \|
	115	+\| CreateStore \| Re-creates/Replaces the original `tl.store` operation using the new pointer expression generated by `CreateAddpr` and the new mask expression generated by `BuildMask` to complete instruction rewriting. \| The prerequisite steps, such as `RewriteAddPtrOp`, `CreateAddpr`, `RewriteStoreOp`, `BuildMask`, must have been successfully executed. \|
	116	+\| RewriteAtomicRWMOp \| Handles pointer issues in atomic read/write/modify operations (such as `atomic.add` and `atomic.max`). \| Generally, it inherits the same limitations as `RewriteAddPtrOp`. Some special, non-contiguous, or conditional atomic operation modes may not be supported. \|
	117	+\| RewriteAtomicCASOp \| Handles pointer linearization issues in atomic compare-and-swap operations (such as `atomic.cas`) and analyzes the pointer expression to eliminate the integer division and modulo operations by tensorization to meet the addressing requirements of hardware atomic instructions. \| \|
	118	+\| RewriteWhile \| Handles pointer overlapping operations in the `while` loop body. \| Complex pointer path transformation that contains conditional branches (`if`) in the loop body is not supported. \|
	119	+\| RewriteFor \| Handles pointer overlapping operations in the `for` loop body. \| \|
	120	+
	121	+#### 3.2.2.2 TritonToUnstructured
	122	+
	123	+\| No.\| Pass/Converter \| Description \|
	124	+\|------\|-------------------------------------------\|-----------------------------------------------------------------------------------------------------------------------------------------------------------------------\|
	125	+\| 1 \| discrete-mask-access-conversion \| Analyzes and converts the memory access pattern based on the discrete index mask in Triton (for example, `triton.language.load` with a non-contiguous mask) to prepare for the subsequent expansion of discrete axes into loops. This pass identifies irregular or sparse access patterns that cannot be efficiently processed by the backend hardware.\|
	126	+\| 2 \| triton-to-unstructured \| Converts the tensor operations identified by `discrete-mask-access-conversion` and containing discrete axes into scalar memory access based on explicit scalar loops.\|
	127	+\| 3 \| bubble-up-operation \| Bubbles up `extract op/extract_slice` for optimization. This can optimize data locality. In some scenarios, unnecessary loops generated after the transformation can be eliminated, thereby improving the execution efficiency of the generated code.\|
	128	+
	129	+
	130	+##### 3.2.2.2.1 discrete-mask-access-conversion
	131	+
	132	+\| Converter \| Description\|
	133	+\|----------------------------\|--------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------\|
	134	+\| DiscreteMaskStoreConversion \| Analyzes the mask and converts the original store operation into the following sequences if it is non-contiguous:<br>1. load (loading the content at the target address)<br>2. select (selecting the target content and the value to be stored based on the mask)<br>3. store (storing the select result back to the target address)\|
	135	+\| DiscreteMaskLoadConversion \| Analyzes the mask and converts the original load operation into the following sequences if it is non-contiguous:<br>1. load (loading all content of the source tensor)<br>2. select (selecting the source tensor content based on the mask, and setting the masked part to the other value) \|
	136	+\| DiscreteMaskAtomicAddConversion \| Analyzes the mask and converts the original atomic_add operation into the following sequences if it is non-contiguous:<br>1. select (selecting the value based on the mask and setting the masked part to 0)<br>2. atomic_add (regenerating the atomic_add operation using the select result)\|
	137	+
	138	+##### 3.2.2.2.2 triton-to-unstructured
	139	+
	140	+\| TritonToUnstructured Converter\| Description\|
	141	+\|---\|---\|
	142	+\| UnstructuredMemAccessConverter\<triton::LoadOp\> \| Converts LoadOp into a multi-loop scalar load operation.\|
	143	+\| UnstructuredMemAccessConverter\<triton::StoreOp\> \| Converts StoreOp into a multi-loop scalar store operation.\|
	144	+\| UnstructuredMemAccessConverter\<triton::AtomicRMWOp\> \| Converts AtomicRMWOp into a multi-loop scalar atomic operation.\|
	145	+\| UnstructuredMemAccessConverter\<triton::AtomicCASOp\> \| Converts AtomicCASOp into a multi-loop scalar atomic operation.\|
	146	+
	147	+##### 3.2.2.2.3 bubble-up-operation
	148	+
	149	+
	150	+\| Converter\| Description\|
	151	+\|---\|---\|
	152	+\| BubbleUpExtract\<tensor::ExtractOp\> \| Bubbles up extract op, avoiding unnecessary loops in some scenarios.\|
	153	+\| BubbleUpExtract\<tensor::ExtractSliceOp\> \| Bubbles up extract op/extract_slice, avoiding unnecessary loops in some scenarios.\|
	154	+
	155	+#### 3.2.2.3 TritonToLinalg
	156	+
	157	+##### 3.2.2.3.1 triton-to-linalg
	158	+
	159	+TritonToLinalg converts ttir to linalg ir.
	160	+
	161	+\| Converter \| Description \|
	162	+\| ------------------------------------------ \| ------------------------------------------------------------ \|
	163	+\| StoreConverter \| triton::StoreOp to memref::copy \|
	164	+\| AddPtrConverter \| triton::AddPtrOp to memref::ReinterpretCast \|
	165	+\| GetProgramIDConverter \| triton::GetProgramIdOp to a param in functionOp \|
	166	+\| GetNumProgramsConverter \| triton::GetNumProgramsOp to a param in functionOp \|
	167	+\| LoadConverter \| triton::LoadOp to memref::copy and bufferization::ToTensorOp \|
	168	+\| AtomicRMWConverter \| triton::AtomicRMWOp to linalg::GenericOp \|
	169	+\| AtomicCASConverter \| triton::AtomicCASOp to linalg::GenericOp \|
	170	+\| MakeRangeConverter \| triton::MakeRangeOp to linalg::GenericOp \|
	171	+\| SplatConverter \| triton::SplatOp to linalg::FillOp \|
	172	+\| ClampFConverter \| triton::ClampFOp to tensor::EmptyOp, linalg::FillOp \|
	173	+\| PreciseDivConverter \| triton::PreciseDivFOp to arith::DivFOp \|
	174	+\| ArgMinConverter \| triton::ArgMinOp to linalg::ReduceOp \|
	175	+\| ArgMaxConverter \| triton::ArgMaxOp to linalg::ReduceOp \|
	176	+\| ReduceConverter \| triton::ReduceOp to linalg::ReduceOp \|
	177	+\| ScanConverter \| triton::ScanOp to func::CallOp \|
	178	+\| ReshapeConverter \| triton::ReshapeOp to tensor::ReshapeOp \|
	179	+\| ExpandDimsConverter \| triton::ExpandDimsOp to tensor::ExpandShapeOp \|
	180	+\| BroadcastConverter \| triton::BroadcastOp to linalg::BroadcastOp \|
	181	+\| DenseConstantConverter \| arith::ConstantOp to linalg::FillOp \|
	182	+\| ExternElementwiseClOpConverter \| triton::ExternElementwiseOp to linalg::MapOp \|
	183	+\| TritonMulhiuiConverter \| triton::MulhiUIOp to arith::MulSIExtendedOp \|
	184	+\| TritonPreciseSqrtConverter \| triton::PreciseSqrtOp to math::SqrtOp \|
	185	+\| AdvanceConverter \| triton::AdvanceOp to memref::ReinterpretCastOp \|
	186	+\| TransposeConverter \| triton::TransOp to linalg::TransposeOp \|
	187	+\| SplitConverter \| triton::SplitOp to tensor::ExtractSliceOp \|
	188	+\| JoinConverter \| triton::JoinOp to tensor::InsertSliceOp \|
	189	+\| CatConverter \| triton::CatOp to tensor::InsertSliceOp \|
	190	+\| BitcastConverter \| triton::BitcastOp to arith::BitcastOp \|
	191	+\| LoopConverter\<scf::ForOp\> \| scf::ForOp to scf::ForOp \|
	192	+\| LoopConverter\<scf::WhileOp\> \| scf::WhileOp to scf::WhileOp \|
	193	+\| YieldConverter \| scf::YieldOp to scf::YieldOp \|
	194	+\| GatherConverter \| triton::GatherOp to func::FuncOp \|
	195	+\| GatherLoadConverter \| triton::GatherLoadOp to scf::ForOp \|
	196	+\| DeviceAssertConverter \| triton::AssertOp to func::FuncOp \|
	197	+\| DevicePrintConverter \| triton::PrintOp to func::FuncOp \|
	198	+\| MatmulConverter \| triton::DotOp to linalg::MatmulOp \|
	199	+\| SortOpConverter \| triton::SortOp to func::FuncOp \|
	200	+\| DotScaledConverter \| triton::DotScaledOp to linalg::MatmulOp \|
	201	+\| PtrToIntConverter \| triton::PtrToIntOp \|
	202	+\| MakeTensorPtrConverter \| triton::PtrToIntOp to arith::IndexCastOp \|
	203	+
	204	+#### 3.2.2.4 other passes
	205	+
	206	+\| Pass\| Description\| Core Converter\| Description\|
	207	+\|---\|---\|---\|---\|
	208	+\| triton-to-annotation \| Converts Ascend NPU-specific compilation hint (`tl.compile_hint`) into backend annotation dialects, which are used to guide subsequent hardware-specific optimization or resource configuration.\| TritonAnnotationConversion \| Converts `triton::AnnotationOp` into `annotation::MarkOp` to transfer advanced compilation hints to the underlying annotation marks.\|
	209	+\| triton-to-hfusion \| Converts `TTIR` in Triton into the corresponding operation in the `HFusion` dialect of the Ascend NPU hardware accelerator.\| TritonHistogramToHFusionConversion \| Converts `triton::HistogramOp` into `hfusion::HistogramOp` to enable efficient execution on the dedicated NPU hardware.\|
	210	+\| triton-to-hivm \| Processes the block synchronization operations (`tl.sync_block_all`, `tl.sync_block_set`, and `tl.sync_block_wait`) of Triton and converts them into the cross-core synchronization instruction in the `HIVM` dialect of Ascend NPU. These instructions are used to manage synchronization and data dependencies in the multi-core pipeline, which is the key to pipeline optimization.\| TritonCustomOpToHIVMSyncOpConversion \| Converts Triton synchronization instructions to HIVM synchronization instructions.<br>• `sync_block_all`: synchronizes blocks globally.<br>• `sync_block_set`: sets a synchronization point.<br>• `sync_block_wait`: waits for a synchronization point.\|
	211	+\| triton-to-llvm \| Converts the inline assembly operation (`tl.inline_assembly`) in Triton to the inline assembly in the LLVM dialect, and finally maps it to a CCE hardware intrinsic function of Ascend NPU.\| ElementwiseInlineAsmOpConversion \| Converts `triton::ElementwiseInlineAsmOp` to `LLVM::InlineAsmOp`.\|
	212	+
	213	+### 3.2.3 Ascend affinitive Operators
	214	+
	215	+\| No.\| Operator \| Description\|
	216	+\|---\|---\|---\|
	217	+\| 1 \| tl.custom_op \| A set of custom operators extended by Ascend NPU, used to support hardware-specific memory access and data movement patterns. For example:<br>• `index_select`: selects data based on an index.<br>• `index_put`: places data based on an index.<br>• `gather_out_to_ub`: collects external data to the unified buffer (UB).<br>• `scatter_ub_to_out`: scatters data from the UB to the output.<br>• `indirect_load`: loads content from an indirect address.<br>• `indirect_store`: stores content to an indirect address.\|
	218	+\| 2 \| tl.compile_hint \| Provides hardware-specific compilation hints to the compiler, which are used to guide the backend optimization policy, resource allocation, or kernel configuration.\|
	219	+\| 3 \| tl.sync_block_wait(`sender, receiver, event_id`) \| Waits for block synchronization. The `receiver` waits for the event signal (`event_id`) sent by the `sender`, which is used to manage data dependencies and execution sequence in the cross-core pipeline.\|
	220	+\| 4 \| tl.sync_block_set(`sender, receiver, event_id`) \| Sets block synchronization. The `sender` sends an event signal (`event_id`) to the `receiver`, indicating that an execution phase or data is ready.\|
	221	+\| 5 \| tl.sync_block_all(`mode, event_id`) \| Globally synchronizes blocks. The sender broadcasts an event signal (`event_id`) to all related receivers according to the specified synchronization mode (`mode`) to implement full-core synchronization or collective synchronization in a specific mode.\|

Adocs/en/debug_guide/debugging.md+503-0

@@ -0,0 +1,503 @@
	1	+# Triton-Ascend Debugging Guide
	2	+
	3	+## 1 Overview
	4	+
	5	+This document is the Triton-Ascend Debugging Guide, which is intended for engineers who participate in adapting Triton to Ascend NPU. It systematically describes the common debugging methods and tools used during Triton-Ascend compilation and running.
	6	+
	7	+The contents of this document are as follows:
	8	+
	9	+\| Section\| Description\|
	10	+\|------\|--------\|
	11	+\| 1. Overview\| Describes the core objectives of debugging (focusing on the `ttir.mlir` → `ttadapter.mlir` conversion) and provides guidance on common issues.\|
	12	+\| 2. Compilation Process Overview\| Describes the key phases of the Triton-Ascend end-to-end compilation chain, providing a context basis for subsequent debugging.\|
	13	+\| 3. Temporary File Guide\| Describes the storage locations and functions of intermediate files (such as the `.mlir`, `.ll`, and`.o` files) generated during the compilation, facilitating manual check.\|
	14	+\| 4. Interpreter Mode\| Describes how to set `TRITON_INTERPRET` to `1` to run the kernel on the CPU and use the result as the accuracy benchmark of the NPU computing result.\|
	15	+\| 5. Debugging Methods\| The following practical debugging methods are provided:<br>• Static/Runtime printing<br>• Compilation error debugging<br>\|
	16	+\| Appendix A\| Provides a quick reference table of common environment variables to improve debugging efficiency.\|
	17	+
	18	+You are advised to refer to the corresponding sections as required to efficiently locate and resolve various exceptions in Triton-Ascend integration.
	19	+
	20	+
	21	+### 1.1 Triton-Ascend Common Issue Classification and Debugging Guide
	22	+
	23	+During development, issues can be classified into different types. The following table provides guidance for quickly identifying issue types and preferred debugging methods.
	24	+
	25	+\| Issue Type\| Typical Symptom/Description\| Preferred Debugging Method\|
	26	+\| :--- \| :--- \| :--- \|
	27	+\| Accuracy issue\| The NPU running result is different from the benchmark reference result (such as the PyTorch or Triton CPU interpreter).\| 4. Interpreter mode<br> 5.1 Debugging by printing\|
	28	+\| Compilation error (MLIRCompileError)\| If the compilation fails in the conversion phase, `MLIRCompileError` is thrown on the Python side.\| 5.2 Compilation error debugging\|
	29	+
	30	+
	31	+## 2 Triton-Ascend Compilation Process Overview
	32	+
	33	+Understanding the complete compilation chain is the basis for effective debugging. The compilation process of Triton-Ascend consists of the following phases:
	34	+
	35	+\| Phase\| Input\| Output\| Tool/Component\| Description\|
	36	+\| :--- \| :--- \| :--- \| :--- \| :--- \|
	37	+\| Python Kernel compilation\| `triton_kernel.py` (Python) \| `ttir.mlir` (MLIR) \| Triton JIT compiler\| Compiles the Triton Python kernel written by users into the standard Triton IR (TTIR).\|
	38	+\| Triton IR adaptation and transformation\| `ttir.mlir` \| `ttadapter.mlir` \| Ascend-adapted Triton backend\| Key debugging phase. Converts TTIR into the adapter IR for the Ascend NPU backend.\|
	39	+\| MLIR compilation and code generation\| `ttadapter.mlir` \| `.o` (executable object file)\| BiSheng compiler (`bishengir-compile`)\| The adapter IR is further compiled and optimized to generate binary code that can be executed on the NPU.\|
	40	+
	41	+```bash
	42	+# Triton-Ascend compilation process
	43	+[Python Kernel]
	44	+ ↓ (triton.compile)
	45	+[ttir.mlir]
	46	+ ↓ │ (TRITON_DEBUG=1 → ~/.triton/dump/)
	47	+[ttadapter.mlir]
	48	+ ↓ (bishengir-compile)
	49	+[NPU executable file.o]
	50	+```
	51	+This guide focuses on the second phase, that is, the `ttir.mlir` → `ttadapter.mlir` conversion. This phase is the main function of Triton-Ascend.
	52	+
	53	+## 3 Triton-Ascend Temporary File Guide
	54	+During the compilation of Triton-Ascend, the system generates multiple temporary files for caching and debugging. Understanding the location and usage of these files is critical for efficient debugging.
	55	+
	56	+### 3.1 Cache
	57	+
	58	+Triton uses the cache mechanism to accelerate the repeated compilation process. Intermediate files generated during compilation are cached in the user directory to avoid repeated compilation of the same kernel.
	59	+
	60	+Cache directory structure:
	61	+
	62	+- Default path: ~/.triton/cache/
	63	+
	64	+Main cache content:
	65	+
	66	+- Input file cache: ttir.mlir file generated by the original Triton kernel
	67	+
	68	+- Output file cache: ttadapter.mlir file converted to adapt to Ascend
	69	+
	70	+- Compilation product cache: executable file generated after compilation
	71	+
	72	+Naming conventions of cache files:
	73	+Cache files are usually named using MD5 hash values to ensure that the same kernel code corresponds to the same cache file.
	74	+
	75	+Recommendations for cache management:
	76	+
	77	+Periodic clearing: Cache files may occupy a large amount of disk space. You can periodically clear the cache files.
	78	+
	79	+```bash
	80	+rm -rf ~/.triton/cache/*
	81	+```
	82	+Disabling cache during debugging: You are advised to temporarily disable the cache to ensure that the compilation is performed each time when debugging compilation issues.
	83	+
	84	+```bash
	85	+export TRITON_DISABLE_CACHE=1
	86	+```
	87	+Cache verification: If you suspect that the issue is caused by the cache, delete related cache files and perform the test again.
	88	+
	89	+### 3.2 Dump Files
	90	+
	91	+You can set the environment variable TRITON_DEBUG to 1 to dump intermediate representation files to disks during compilation. These files are key resources for debugging compilation issues.
	92	+
	93	+Dump directory structure:
	94	+
	95	+- Default path: ~/.triton/dump/
	96	+
	97	+Directory naming: A subdirectory named by a timestamp or unique ID is generated for each compilation session.
	98	+
	99	+Main dump files:
	100	+
	101	+- kernel.ttir.mlir: Triton IR file (compilation input)
	102	+
	103	+- kernel.ttadapter.mlir: adapter IR file (conversion output)
	104	+
	105	+Enabling debug dump:
	106	+Even if the cache is enabled, the system still generates dump files (overriding files in the directory with the same name) each time the system runs as long as TRITON_DEBUG=1 is set. However, if the cache is hit and compilation is skipped, IR conversion may not be triggered. As a result, no new dump file is generated. Therefore, during debugging, you are advised to set as follows:
	107	+```bash
	108	+# Set environment variables before running the Triton program.
	109	+export TRITON_DEBUG=1
	110	+export TRITON_DISABLE_CACHE=1
	111	+
	112	+# Run Triton kernel.
	113	+python your_triton_program.py
	114	+```
	115	+
	116	+### 3.3 File Lifecycle Management
	117	+Understanding when these temporary files are generated and how they are cleared helps you better manage the debugging environment.
	118	+
	119	+File generation time table
	120	+
	121	+\| File Type\| Generation Phase\| Triggering Condition\| Clearance Suggestion\|
	122	+\|----------\|----------\|----------\|----------\|
	123	+\| Cache file\| During each compilation\| Generated when the cache is not hit\| Periodic clearing or clearing during troubleshooting\|
	124	+\| Dump file\| After TRITON_DEBUG=1 is set\| Generated during each compilation\| Manual clearing after debugging\|
	125	+
	126	+- In the production environment, debug dump should be disabled (that is, TRITON_DEBUG=1 is not set).
	127	+
	128	+- The cache mechanism can significantly improve performance and should not be disabled.
	129	+
	130	+By properly using these temporary files, developers can efficiently locate and solve issues encountered during Triton-Ascend compilation.
	131	+
	132	+### 3.4 IR File Parsing
	133	+
	134	+The following uses the [01-vector-add.py](../../../third_party/ascend/tutorials/01-vector-add.py#) test case as an example to describe the compilation process:
	135	+This is a simple addition calculation of two tensors. For the calculation logic, see the comments in the sample case.
	136	+You can enable the dump file output by setting TRITON_DEBUG=1 to obtain kernel.ttir.mlir and kernel.ttadapter.mlir.
	137	+- Run the test case.
	138	+```
	139	+TRITON_DEBUG=1 python 01-vector-add.py
	140	+```
	141	+After the test case is executed, the dump file path is displayed. The default path is ~/.triton/dump. The following information is displayed:
	142	+```
	143	+Dumping intermediate results to ~/.triton/dump/xxx
	144	+# xxx is a unique hash identifier.
	145	+```
	146	+Go to the dump path and view kernel.ttir.mlir and kernel.ttadapter.mlir.
	147	+
	148	+#### 3.4.1 Triton Intermediate Representation (TTIR)
	149	+- TTIR example
	150	+The kernel.ttir.mlir file is as follows:
	151	+```
	152	+module {
	153	+ tt.func public @add_kernel(%arg0: !tt.ptr<f32> {tt.divisibility = 16 : i32} , %arg1: !tt.ptr<f32> {tt.divisibility = 16 : i32} , %arg2: !tt.ptr<f32> {tt.divisibility = 16 : i32} , %arg3: i32 {tt.divisibility = 16 : i32} ) attributes {noinline = false} {
	154	+ %cst = arith.constant dense<0.000000e+00> : tensor<1024xf32> loc(#loc1)
	155	+ %c1024_i32 = arith.constant 1024 : i32 loc(#loc1)
	156	+ %0 = tt.get_program_id x : i32 loc(#loc2)
	157	+ %1 = arith.muli %0, %c1024_i32 : i32 loc(#loc3)
	158	+ %2 = tt.make_range {end = 1024 : i32, start = 0 : i32} : tensor<1024xi32> loc(#loc4)
	159	+ %3 = tt.splat %1 : i32 -> tensor<1024xi32> loc(#loc5)
	160	+ %4 = arith.addi %3, %2 : tensor<1024xi32> loc(#loc5)
	161	+ %5 = tt.splat %arg3 : i32 -> tensor<1024xi32> loc(#loc6)
	162	+ %6 = arith.cmpi slt, %4, %5 : tensor<1024xi32> loc(#loc6)
	163	+ %7 = tt.splat %arg0 : !tt.ptr<f32> -> tensor<1024x!tt.ptr<f32>> loc(#loc7)
	164	+ %8 = tt.addptr %7, %4 : tensor<1024x!tt.ptr<f32>>, tensor<1024xi32> loc(#loc7)
	165	+ %9 = tt.load %8, %6, %cst : tensor<1024x!tt.ptr<f32>> loc(#loc8)
	166	+ %10 = tt.splat %arg1 : !tt.ptr<f32> -> tensor<1024x!tt.ptr<f32>> loc(#loc9)
	167	+ %11 = tt.addptr %10, %4 : tensor<1024x!tt.ptr<f32>>, tensor<1024xi32> loc(#loc9)
	168	+ %12 = tt.load %11, %6, %cst : tensor<1024x!tt.ptr<f32>> loc(#loc10)
	169	+ %13 = arith.addf %9, %12 : tensor<1024xf32> loc(#loc11)
	170	+ %14 = tt.splat %arg2 : !tt.ptr<f32> -> tensor<1024x!tt.ptr<f32>> loc(#loc12)
	171	+ %15 = tt.addptr %14, %4 : tensor<1024x!tt.ptr<f32>>, tensor<1024xi32> loc(#loc12)
	172	+ tt.store %15, %13, %6 : tensor<1024x!tt.ptr<f32>> loc(#loc13)
	173	+ tt.return loc(#loc14)}}
	174	+```
	175	+- TTIR analysis
	176	+
	177	+TTIR is an intermediate representation generated by the frontend of the Triton compiler. It is expressed in the Multi-Level IR (MLIR) format and retains the semantic structure of the original Triton Python kernel. In `kernel.ttir.mlir`:
	178	+
	179	+- The `@add_kernel` function receives three pointer parameters (corresponding to the device memory addresses of input A, input B, and output C respectively) and an integer parameter `n` indicating the vector length.
	180	+- Each triton program (vectorized execution unit) processes 1024 elements (represented by the `%c1024_i32` constant), obtains the ID of the current block through the `tt.get_program_id x`, and calculates the global offset.
	181	+- `tt.make_range` and `tt.splat` are used to construct a SIMD-style index tensor and they are used together with `arith.addi` to generate the global address offset processed by each thread.
	182	+- `tt.addptr` and `tt.load` are used to implement vectorized loading, and the mask `%6` (generated by `arith.cmpi slt`) is used to prevent out-of-bounds access.
	183	+- The element-wise floating-point addition `arith.addf` is executed, and the result is returned to the global memory by using `tt.store`.
	184	+
	185	+The TTIR layer is still based on the native abstraction (such as `!tt.ptr<f32>`, `tt.load`, and `tt.store`) of Triton and has not been mapped to the specific memory model or execution unit of the underlying hardware. It is a platform-independent high-level IR.
	186	+
	187	+#### 3.4.1 Target-Specific Adapter Representation (TTAdapter IR)
	188	+
	189	+- TTAdapter IR example
	190	+The kernel.ttadapter.mlir file is as follows:
	191	+```
	192	+module {
	193	+ func.func @add_kernel(%arg0: memref<?xi8>, %arg1: memref<?xi8>, %arg2: memref<?xf32> {tt.divisibility = 16 : i32, tt.tensor_kind = 0 : i32}, %arg3: memref<?xf32> {tt.divisibility = 16 : i32, tt.tensor_kind = 0 : i32}, %arg4: memref<?xf32> {tt.divisibility = 16 : i32, tt.tensor_kind = 1 : i32}, %arg5: i32 {tt.divisibility = 16 : i32}, %arg6: i32, %arg7: i32, %arg8: i32, %arg9: i32, %arg10: i32, %arg11: i32) attributes {SyncBlockLockArgIdx = 0 : i64, WorkspaceArgIdx = 1 : i64, global_kernel = "local", mix_mode = "aiv", parallel_mode = "simd"} {
	194	+ %cst = arith.constant 0.000000e+00 : f32
	195	+ %c1024 = arith.constant 1024 : index
	196	+ %c1024_i32 = arith.constant 1024 : i32
	197	+ %0 = arith.muli %arg9, %c1024_i32 : i32
	198	+ %1 = arith.index_cast %0 : i32 to index
	199	+ %reinterpret_cast = memref.reinterpret_cast %arg2 to offset: [%1], sizes: [1024], strides: [1] : memref<?xf32> to memref<1024xf32, strided<[1], offset: ?>>
	200	+ %alloc = memref.alloc() : memref<1024xf32>
	201	+ %2 = arith.addi %1, %c1024 : index
	202	+ %3 = arith.index_cast %arg5 : i32 to index
	203	+ %4 = arith.maxsi %1, %3 : index
	204	+ %5 = arith.minsi %2, %4 : index
	205	+ %6 = arith.subi %5, %1 : index
	206	+ %7 = arith.cmpi slt, %6, %c1024 : index
	207	+ scf.if %7 {
	208	+ linalg.fill ins(%cst : f32) outs(%alloc : memref<1024xf32>)
	209	+ } {hivm.unlikely_condition}
	210	+ %subview = memref.subview %reinterpret_cast[0] [%6] [1] : memref<1024xf32, strided<[1], offset: ?>> to memref<?xf32, strided<[1], offset: ?>>
	211	+ %subview_0 = memref.subview %alloc[0] [%6] [1] : memref<1024xf32> to memref<?xf32, strided<[1]>>
	212	+ memref.copy %subview, %subview_0 : memref<?xf32, strided<[1], offset: ?>> to memref<?xf32, strided<[1]>>
	213	+ %8 = bufferization.to_tensor %alloc restrict writable : memref<1024xf32>
	214	+ %reinterpret_cast_1 = memref.reinterpret_cast %arg3 to offset: [%1], sizes: [1024], strides: [1] : memref<?xf32> to memref<1024xf32, strided<[1], offset: ?>>
	215	+ %alloc_2 = memref.alloc() : memref<1024xf32>
	216	+ scf.if %7 {
	217	+ linalg.fill ins(%cst : f32) outs(%alloc_2 : memref<1024xf32>)
	218	+ } {hivm.unlikely_condition}
	219	+ %subview_3 = memref.subview %reinterpret_cast_1[0] [%6] [1] : memref<1024xf32, strided<[1], offset: ?>> to memref<?xf32, strided<[1], offset: ?>>
	220	+ %subview_4 = memref.subview %alloc_2[0] [%6] [1] : memref<1024xf32> to memref<?xf32, strided<[1]>>
	221	+ memref.copy %subview_3, %subview_4 : memref<?xf32, strided<[1], offset: ?>> to memref<?xf32, strided<[1]>>
	222	+ %9 = bufferization.to_tensor %alloc_2 restrict writable : memref<1024xf32>
	223	+ %10 = arith.addf %8, %9 : tensor<1024xf32>
	224	+ %reinterpret_cast_5 = memref.reinterpret_cast %arg4 to offset: [%1], sizes: [1024], strides: [1] : memref<?xf32> to memref<1024xf32, strided<[1], offset: ?>>
	225	+ %extracted_slice = tensor.extract_slice %10[0] [%6] [1] : tensor<1024xf32> to tensor<?xf32>
	226	+ %subview_6 = memref.subview %reinterpret_cast_5[0] [%6] [1] : memref<1024xf32, strided<[1], offset: ?>> to memref<?xf32, strided<[1], offset: ?>>
	227	+ bufferization.materialize_in_destination %extracted_slice in writable %subview_6 : (tensor<?xf32>, memref<?xf32, strided<[1], offset: ?>>) -> ()
	228	+ return
	229	+ }
	230	+}
	231	+```
	232	+
	233	+- TTAdapter IR parsing
	234	+
	235	+TTIR is converted to TTAdapter IR to adapt to the Ascend NPU architecture in the Triton-Ascend compilation process. TTAdapter IR uses standard MLIR dialect (such as `memref`, `linalg`, and `scf`) and introduces NPU-specific constraints and optimization policies. In `kernel.ttadapter.mlir`:
	236	+
	237	+- The function signature has been converted from the Triton pointer type to `memref<?xi8>` or `memref<?xf32>` with attributes. `tt.divisibility = 16` indicates the memory alignment requirement, and `tt.tensor_kind` distinguishes input (marked with 0) and output (marked with 1).
	238	+- The global offset is reconstructed as a local view of a fixed size (1024) by using `memref.reinterpret_cast` for subsequent vectorization.
	239	+- The boundary check logic is introduced to calculate the number of valid elements `%6` and use `scf.if` to control whether to fill zeros (`linalg.fill`) at the end to ensure that the SIMD width is aligned and does not exceed the boundary.
	240	+- `memref.alloc` is used to allocate a local buffer, `memref.copy` is used to securely copy the global memory data to the local host, and `bufferization.to_tensor` is used to convert the data into tensors for operators.
	241	+- The addition operation is performed by `arith.addf` on the tensor. The valid part of the result is truncated by `tensor.extract_slice` and written back to the target memref by `bufferization.materialize_in_destination`.
	242	+
	243	+TTAdapter IR has been abstracted from Triton to adapt to the Ascend NPU format.
	244	+
	245	+## 4 Interpreter Mode
	246	+The core value of the interpreter is to isolate hardware differences. You can set the environment variable `TRITON_INTERPRET` to `1` to forcibly execute kernel computation on the CPU. The result of the kernel computation can be used as the benchmark for determining the NPU computation accuracy.
	247	+
	248	+Usage:
	249	+1. Set the environment variable `TRITON_INTERPRET` to `1` and run the program so that the Triton kernel is executed on the CPU interpreter.
	250	+2. Insert a Python breakpoint at the position to be checked in the Triton kernel source code.
	251	+ ```python
	252	+ breakpoint() # Python built-in breakpoint function
	253	+ ```
	254	+3. The program execution is paused and you enter the Python debugger (`Pdb`). You can print and check the value of any intermediate variable.
	255	+ ```python
	256	+ (Pdb) p tmp0 # Print the value of variable tmp0.
	257	+ ```
	258	+
	259	+- Note: The interpreter mode performs all computations on the CPU, which significantly reduces the running efficiency. Therefore, after debugging or verification, you must cancel the setting of the environment variable TRITON_INTERPRET or explicitly set it to 0 to ensure that the system performance is not affected.
	260	+
	261	+```bash
	262	+# Cancel the environment variable.
	263	+unset TRITON_INTERPRET
	264	+
	265	+# Explicitly set it to 0.
	266	+export TRITON_INTERPRET=0
	267	+```
	268	+
	269	+## 5 Debugging Methods
	270	+
	271	+### 5.1 Debugging by Printing
	272	+### 5.1.1 Static Printing Debugging
	273	+This method uses `tl.static_print` to print the value of a constant expression during compilation. It is applicable to debugging configuration parameters and constants that are known during compilation.
	274	+
	275	+Setting the environment variable `TRITON_DEVICE_PRINT` to `1` can enable the `tl.static_print` function. This function allows constant values to be printed during kernel compilation. It is an effective method for verifying configuration parameters and constant expressions.
	276	+
	277	+Features:
	278	+
	279	+- `tl.static_print` is executed during compilation, not during runtime.
	280	+
	281	+- Only compilation constants (tl.constexpr parameters and constant expressions) can be printed.
	282	+
	283	+- The output is displayed in the standard output of the compiler.
	284	+
	285	+Usage:
	286	+
	287	+1. In the Triton kernel, add the `tl.static_print` statement for the constant parameters to be debugged.
	288	+```python
	289	+import triton.language as tl
	290	+
	291	+@triton.jit
	292	+def triton_kernel(
	293	+ out_ptr0,
	294	+ in_ptr0,
	295	+ in_ptr1,
	296	+ XBLOCK: tl.constexpr, # Constant parameter during compilation
	297	+ USE_FP16: tl.constexpr # Constant parameter during compilation
	298	+):
	299	+ # Print constant parameters during compilation.
	300	+ tl.static_print("XBLOCK = ", XBLOCK)
	301	+ tl.static_print("USE_FP16 = ", USE_FP16)
	302	+
	303	+ idx = tl.arange(0, XBLOCK)
	304	+ tmp0 = tl.load(in_ptr0 + idx)
	305	+ tmp1 = tl.load(in_ptr1 + idx)
	306	+
	307	+ # Print the constant calculation result.
	308	+ elements_per_thread = XBLOCK // 32
	309	+ tl.static_print("Elements per thread = ", elements_per_thread)
	310	+
	311	+ tmp2 = tmp0 + tmp1
	312	+ tl.store(out_ptr0 + idx, tmp2)
	313	+```
	314	+2. Set the environment variable and run the program for compilation.
	315	+```bash
	316	+# Enable Triton debugging output (including static_print).
	317	+export TRITON_DEVICE_PRINT=1
	318	+
	319	+# Run the Python program. The output is displayed in the compilation phase.
	320	+python your_program.py
	321	+```
	322	+
	323	+
	324	+### 5.1.2 Runtime Debugging
	325	+You can use `tl.device_print` to flexibly print the values of the variables to be observed.
	326	+Setting the environment variable `TRITON_DEVICE_PRINT` to `1` can enable the `tl.device_print` function. This function allows tensor values to be printed in the kernel. It is an efficient method for verifying the computation accuracy by phase.
	327	+
	328	+Usage:
	329	+1. In the Triton kernel, add the `tl.device_print` statement for the variables to be printed.
	330	+```python
	331	+import triton.language as tl
	332	+
	333	+@triton.jit
	334	+def triton_kernel(out_ptr0, in_ptr0, in_ptr1, XBLOCK: tl.constexpr):
	335	+ idx = tl.arange(0, XBLOCK)
	336	+ tmp0 = tl.load(in_ptr0 + idx)
	337	+ tmp1 = tl.load(in_ptr1 + idx)
	338	+ tmp2 = tmp0 + tmp1
	339	+ tl.device_print("tmp2 after addition = ", tmp2) # Print the intermediate result.
	340	+ tl.store(out_ptr0 + idx, tmp2)
	341	+```
	342	+2. Set the environment variable `TRITON_DEVICE_PRINT` to `1` and run the program. The window displays the value of the variable.
	343	+```bash
	344	+# Enable Triton debugging output (including device_print).
	345	+export TRITON_DEVICE_PRINT=1
	346	+
	347	+# Run the Python program. The output is displayed in the compilation phase.
	348	+python your_program.py
	349	+```
	350	+
	351	+- Note: The print length is limited.
	352	+The length of the tensor printed by `tl.device_print` is limited. When the tensor length exceeds a certain threshold, the output is truncated.
	353	+
	354	+### 5.1.3 Comparing the Two Printing Methods
	355	+
	356	+\| Feature\| `tl.device_print` \| `tl.static_print` \|
	357	+\|------\|-------------------\|-------------------\|
	358	+\| Execution time\| Runtime (kernel execution)\| Compilation (kernel compilation)\|
	359	+\| Output location\| Runtime standard output\| Compiler standard output\|
	360	+\| Print content\| Runtime tensor values and variables\| Compilation constants and constant expressions\|
	361	+\| Impact on performance\| There is runtime overhead.\| No runtime overhead.\|
	362	+\| Enabling environment variables\| `TRITON_DEVICE_PRINT=1` \| `TRITON_DEVICE_PRINT=1` \|
	363	+
	364	+Description of environment variables:
	365	+
	366	+TRITON_DEVICE_PRINT=1: enables runtime printing and compilation printing.
	367	+
	368	+TRITON_DEBUG=1: enables all debugging outputs (including compilation and runtime printing).
	369	+
	370	+### 5.2 Compilation Error Debugging
	371	+When the `ttir.mlir` → `ttadapter.mlir` conversion fails, the `ttadapter.mlir` cannot be generated and the `MLIRCompilationError` error is reported.
	372	+You need to locate the fault at the Triton-Ascend code layer. Triton-Ascend contains the Python and C++ code layers. You need to locate the error code segment based on the call stack information in the error log and use the corresponding debugging method.
	373	+
	374	+### 5.2.1 Debugging Python Code
	375	+When the call stack information shows that the error is caused by the Python layer code of Triton-Ascend, you can use the built-in debugger pdb of Python for interactive debugging. As an effective tool for locating Python code logic errors, pdb allows you to set breakpoints, perform step-by-step execution, and check variable status.
	376	+
	377	+Procedure:
	378	+
	379	+Locating faults
	380	+In the error log, find the Python call stack information closest to the user code, which is usually near the top of the stack. For example:
	381	+
	382	+```text
	383	+File "/path/to/triton/ascend/compiler.py", line 123, in compile_fn
	384	+ result = lower_function(...)
	385	+```
	386	+
	387	+Inserting a debugging breakpoint
	388	+Insert a pdb breakpoint in the Python source file that is suspected to be faulty.
	389	+
	390	+```python
	391	+def compile_fn(ttir):
	392	+ import pdb; pdb.set_trace() # Compatible with all Python versions
	393	+```
	394	+
	395	+Example:
	396	+Assume that a breakpoint is set in line 123 of `compiler.py`. After the program is suspended, the following information is displayed:
	397	+```python
	398	+python
	399	+(Pdb) l # View the current code context.
	400	+118 def compile_fn(ttir):
	401	+120 import pdb; pdb.set_trace()
	402	+121 # Check the input parameter.
	403	+122 print(f"ttir type: {type(ttir)}")
	404	+123 result = lower_function(ttir) # <-- The current suspension position.
	405	+
	406	+(Pdb) p ttir # Check the input parameter.
	407	+(Pdb) n # Execute the next line of code.
	408	+(Pdb) p result # View the result.
	409	+```
	410	+
	411	+### 5.2.2 Debugging Environment Variables
	412	+
	413	+When developing or debugging Triton operators, you can set the following environment variables to enable IR printing in different phases, which helps locate faults. The following describes the two key debugging switches.
	414	+
	415	+#### 5.2.2.1 `MLIR_ENABLE_DUMP=1`
	416	+
	417	+Function:
	418	+Enables automatic dump of the MLIR high-level IR and outputs the IR of the current function in readable text to `stderr` before and after each MLIR pass is executed.
	419	+
	420	+Feature:
	421	+Small log size: usually dozens to hundreds of lines, which are easy to read.
	422	+Focus on high-level logic: applicable to debugging operator conversion, memory layout, and parallel policies.
	423	+
	424	+Suggestion:
	425	+First choice for routine debugging: This log can be used to locate 90% of Triton operator issues.
	426	+It can be used together with `TRITON_DEBUG=1` to further enhance information.
	427	+
	428	+Enabling method:
	429	+```bash
	430	+export MLIR_ENABLE_DUMP=1
	431	+export TRITON_DEBUG=1
	432	+python your_triton_script.py
	433	+```
	434	+
	435	+#### 5.2.2.2 `TRITON_ENABLE_LLVM_DEBUG=1`
	436	+
	437	+Function:
	438	+Enables full debugging logs in the LLVM backend CodeGen phase, including instruction selection, register allocation, instruction scheduling, and machine code generation.
	439	+
	440	+Feature:
	441	+Large log size: A single kernel can generate tens of thousands of lines of output.
	442	+Bottom-layer details: Register name, physical/virtual register mapping, and stack frame layout are included.
	443	+Only for LLVM experts: For common Triton developers, this is considered "noise."
	444	+
	445	+Suggestion:
	446	+Enable this function only when LLVM backend bugs are suspected (for example, invalid instructions are generated or performance exceptions occur).
	447	+It can be used together with LLVM_DEBUG_ONLY to limit the output scope.
	448	+
	449	+When `TRITON_ENABLE_LLVM_DEBUG=1` is enabled, you can use the `LLVM_DEBUG_ONLY` environment variable to specify the module for which the logs will be output. The following is a brief description of the common `DEBUG_TYPE`:
	450	+
	451	+```bash
	452	+## `isel` (Instruction Selection)
	453	+- Function: Converts LLVM IR instructions into machine instructions (MachineInstr) of the target architecture.
	454	+- Debugging content: Displays the mapping process and pattern matching result between IR and machine instructions.
	455	+- Application scenario: The instruction selection is suspected to be incorrect (for example, invalid instructions or inefficient instruction sequences are generated).
	456	+
	457	+## `regalloc` (Register Allocation)
	458	+- Function: Allocates physical registers to virtual registers and processes spilling.
	459	+- Debugging content: Status before and after register allocation, conflict graph, and active interval analysis.
	460	+- Application scenario: The register pressure is high, the performance deteriorates, or unexpected memory access occurs.
	461	+
	462	+## `spiller` (Spiller)
	463	+- Function: Spills some values to the stack memory when registers are insufficient.
	464	+- Debugging content: Which virtual registers are spilled and the positions of inserted load/store instructions.
	465	+- Application scenario: The performance deteriorates due to frequent memory access, and register usage needs to be optimized.
	466	+
	467	+## `peephole` (Peephole Optimizer)
	468	+- Function: Performs partial optimization (such as constant folding and redundant instruction elimination) at the machine code layer.
	469	+- Debugging content: Comparison of instructions before and after optimization.
	470	+- Application scenario: The generated code is redundant, but high-level optimization is not overriding.
	471	+
	472	+## `asm-printer` (Assembly Printer)
	473	+- Function: Converts MachineInstr into the final assembly text (such as PTX, AMDGCN, and CCE).
	474	+- Debugging content: Generated assembly code, symbol references, and instruction encoding.
	475	+- Application scenario: Assembly syntax errors, tag mismatch, or viewing final output.
	476	+```
	477	+
	478	+Enabling method:
	479	+In the following example, it is specified that only `isel` is output.
	480	+```bash
	481	+export TRITON_ENABLE_LLVM_DEBUG=1
	482	+export LLVM_DEBUG_ONLY="isel"
	483	+python your_triton_script.py
	484	+```
	485	+Recommended debugging process:
	486	+Enable `MLIR_ENABLE_DUMP=1` first.
	487	+→ Check whether the conversion at the MLIR layer is correct (for example, ReduceOp → scf.for).
	488	+If the MLIR is normal but the result is incorrect:
	489	+→ It is suspected that the LLVM is faulty. Enable `TRITON_ENABLE_LLVM_DEBUG=1 + LLVM_DEBUG_ONLY`.
	490	+Do not directly enable `TRITON_ENABLE_LLVM_DEBUG=1`.
	491	+→ Large log size may mask key information and severely affect the running speed.
	492	+
	493	+
	494	+## Appendix A: Quick Reference Table for Common Environment Variables
	495	+
	496	+\| Variable \| Description \|
	497	+\|--------------------------\|----------------------------------\|
	498	+\| `TRITON_DEBUG=1` \| Enables intermediate IR dump. \|
	499	+\| `TRITON_DISABLE_CACHE=1` \| Disables compilation cache. \|
	500	+\| `TRITON_INTERPRET=1` \| Uses the CPU interpreter to execute the kernel. \|
	501	+\| `TRITON_DEVICE_PRINT=1` \| Enables runtime print output and compilation print output. \|
	502	+\| `MLIR_ENABLE_DUMP=1` \| Enables automatic dump of the MLIR high-level IR. Outputs the IR of the current function in readable text before and after each MLIR pass is executed.\|
	503	+\| `TRITON_ENABLE_LLVM_DEBUG=1` \| Enables full debugging logs in the LLVM backend CodeGen phase, including instruction selection, register allocation, instruction scheduling, and machine code generation.\|

Adocs/en/debug_guide/profiling.md+162-0

@@ -0,0 +1,162 @@
	1	+# Triton-Ascend Performance Analysis Method
	2	+
	3	+## Obtaining Performance Data
	4	+Before performance optimization, you need to obtain accurate performance data, understand the current performance status, and analyze the next optimization direction based on the current performance status. MindStudio provides multiple methods for testing the performance of the Triton operator, including board profiling and single-operator performance simulation pipeline.
	5	+
	6	+### Board Profiling
	7	+The msProf performance analysis tool is used to collect and analyze key performance metrics of operators running on Ascend AI Processors. You can efficiently locate software and hardware performance bottlenecks of operators based on the output performance data, thereby enhancing the overall efficiency of operator performance analysis.
	8	+- Note: The msProf tool depends on the msopprof executable file in the CANN package. The interface functions in this file are the same as those in msprof op. This file is provided by the CANN package and does not need to be installed separately. For details about common msProf commands, see [Common msProf Commands](https://www.hiascend.com/document/detail/zh/mindstudio/82RC1/ODtools/Operatordevelopmenttools/atlasopdev_16_0082.html).
	9	+
	10	+The following command is an example of collecting performance data of an operator on a board. You can flexibly combine and configure parameters as required. In the example, --output is an optional parameter for specifying the path for storing the collected performance data. --kernel-name is an optional parameter for specifying the performance data of a single kernel to be collected. If you want to collect the performance data of all operators, you do not need to specify --kernel-name. $HOME/projects/test_op.py is the executable script of the operator.
	11	+```
	12	+msprof op --kernel-name=target_kernel_name --output=$HOME/projects/output python3 $HOME/projects/test_op.py
	13	+```
	14	+The following uses the [03-layer-norm.py](../../../third_party/ascend/tutorials/03-layer-norm.py) test case as an example (the generated data file is saved in the current path when if --output is not specified):
	15	+```
	16	+msprof op --kernel-name=_layer_norm_fwd_fused python3 03-layer-norm.py
	17	+```
	18	+- Note: For details about the result data of all the following collection items, see [op_summary (Operator Details)](https://www.hiascend.com/document/detail/zh/canncommercial/83RC1/devaids/Profiling/atlasprofiling_16_0067.html) in the CANN Performance Optimization Tool User Guide.
	19	+Figure 1 PipeUtilization.csv (ratios of time taken by compute units and MTEs)
	20	+![alt text](../figures/time_consumed.png)
	21	+
	22	+### Operator Simulation Pipeline Diagram
	23	+The operator optimization tool msProf supports profile data collection and automatic parsing in a simulation environment. For details about how to obtain the simulation pipeline diagram by using the msProf tool, see [Pipeline diagram](https://www.hiascend.com/document/detail/zh/canncommercial/83RC1/devaids/optool/atlasopdev_16_0087.html).
	24	+The command for generating the operator simulation pipeline diagram is similar to that for collecting operator board performance data. The preceding `03-layer-norm.py` is used as an example. `--soc-version` is used to specify the hardware version of the current machine. You can enter `npu-smi info` in the terminal to view the hardware version.
	25	+```
	26	+# Path of the source simulator
	27	+export LD_LIBRARY_PATH=/root/CANN/Install_CANN/Ascend/ascend_toolkit/latest/tools/simulator/{soc-version}/lib:$LD_LIBRARY_PATH
	28	+# Collecting the operator simulation pipeline diagram
	29	+msprof op simulator --kernel-name=_layer_norm_fwd_fused --soc-version={soc-version} python3 03-layer-norm.py
	30	+```
	31	+- Note: In the preceding example, `soc-version=Ascend910B3`.
	32	+
	33	+\|Soc-Version\|
	34	+\| :---: \| :---: \| :---: \|
	35	+\|Ascend910A\|Ascend310\|Ascend310B1\|
	36	+\|Ascend910B\|Ascend310P1\|Ascend310B2\|
	37	+\|Ascend910B1\|Ascend310P2\|Ascend310B3\|
	38	+\|Ascend910B2\|Ascend310P3\|Ascend310B4\|
	39	+\|Ascend910B2C\|Ascend310P4\|-\|
	40	+\|Ascend910B3\|Ascend310P5\|-\|
	41	+\|Ascend910B4\|Ascend310P7\|-\|
	42	+
	43	+The following two files save the obtained performance data:
	44	+- trace.json
	45	+- visualize_data.bin
	46	+
	47	+The trace.json file supports the following two visualized display modes:
	48	+- Chrome browser
	49	+ Enter the `chrome://tracing` address in the address box of the Chrome browser, drag the instruction pipeline file (trace.json) generated by msprof op simulator to the blank area, and press the shortcut keys on the keyboard (W: zoom in; S: zoom out; A: move left; D: move right) to view the file.
	50	+ Figure 2 Timeline page on Chrome
	51	+![alt text](../figures/trace_json_with_chrome.png)
	52	+
	53	+- [MindStudio Insight](https://www.hiascend.com/document/detail/zh/mindstudio/82RC1/GUI_baseddevelopmenttool/msascendinsightug/Insight_userguide_0005.html) visualized display
	54	+MindStudio Insight provides the running status of instructions on the Ascend AI Processor in a sequence diagram. Users can identify the sequence optimization points of micro instructions by analyzing the instruction details, instruction execution time, call stack of the code associated with the instruction, and synchronization lines between instructions and pipelines in the sequence diagram.
	55	+ Figure 3 Timeline page on MindStudio Insight
	56	+ ![alt text](../figures/trace_json_with_insight.png)
	57	+
	58	+The visualize_data.bin file can be visualized on MindStudio Insight.
	59	+- In addition to collecting performance data like trace.json, visualize_data.bin also provides an instruction association dashboard corresponding to the source code (for example, 03-layer-norm.py).
	60	+ Figure 4 MindStudio Insight-visualize_data.bin instruction association
	61	+ - Note: For details about the result data of the following collection items, see [Operator Optimization](https://www.hiascend.com/document/detail/zh/mindstudio/82RC1/GUI_baseddevelopmenttool/msascendinsightug/Insight_userguide_0068.html) in MindStudio Insight.
	62	+ ![alt text](../figures/visualize_data_with_insight.png)
	63	+
	64	+
	65	+## Analyzing Performance Data
	66	+
	67	+### Theoretical Parameters
	68	+The theoretical performance is the ideal objective of the actual performance of the operator. Different hardware platforms have different hardware specifications. Theoretical performance helps us understand the potential of hardware and set performance optimization objectives.
	69	+
	70	+- Theoretical time required for transfer-related pipelines (such as MTE1, MTE2, and MTE3) = Data volume (unit: byte)/Theoretical bandwidth. For example, if the peak GM bandwidth of an AI processor is about 1.8 TB/s, the theoretical time required for transferring a float-type matrix of size 4096 x 4096 is sizeof(float) x 4096 x 4096/1.8 TB/s = 37.28 μs (calculated based on 1 TB = 10<sup>12</sup> Byte).
	71	+> Note:
	72	+> - If multiple transfer instructions exist at the same time, the bandwidth is shared. Data cannot be moved at a rate close to the theoretical bandwidth. For example, if the MTE2 and MTE3 read and write the GM at the same time, the time consumed by the transfer pipeline is (MTE2 transfer volume + MTE3 transfer volume)/GM bandwidth.
	73	+> - The bandwidth usage (effective bandwidth/theoretical bandwidth) varies according to the size of data blocks to be transferred. If the amount of data transferred each time is small, the actual performance cannot reach the theoretical bandwidth.
	74	+- Theoretical time required for compute-related pipelines (such as Cube, Vector, and Scalar) = Data volume (unit: element)/Theoretical computing power. For example, if the theoretical peak computing power of a certain AI processor for float data type vectors is 11.06 TOPS, the theoretical time required for performing a single instruction computation of 32K float elements is 32K/11.06 TOPS = 0.003 μs (calculated based on 1K = 1000).
	75	+
	76	+### Locating Bottlenecks
	77	+After the performance data is obtained, processes that deviate significantly from theoretical values or consume excessive time are identified as "bottlenecks." The following describes how to find bottlenecks and corresponding optimization directions based on performance data.
	78	+
	79	+- Method 1: Use board profiling to analyze the pipeline.
	80	+View the *op_summary_\.csv** file parsed by board profiling to analyze the pipeline. Note that "\*" indicates the timestamp.
	81	+![analyse_data_op_summary](../figures/performance_analysis_analyse_data_op_summary.png)
	82	+
	83	+ In ideal cases, the utilization rate of each pipeline should be 100%. Any pipeline falling short of this target represents room for improvement. The preceding figure shows the data obtained from an AI processor. In the first scenario of the Vector operator _layer_norm_fwd_fused, the Vector pipeline utilization aiv_vec_ratio is less than 10%, indicating that the computing power is not fully utilized. The Scalar pipeline utilization aiv_scalar_ratio is about 60%, indicating that Scalar is the longest pipeline. \
	84	+ When Scalar is the longest pipeline, analyze whether complex operations are performed on scalar values in the operator source code. The SIMD microarchitecture of Ascend is more suitable for multi-data parallel computing. Another possibility is that the Triton software stack degrades vector computing to scalar computing because some instructions do not support specific data types on the hardware. Optimization should involve both pipeline and scalar optimization methods. For details, see method 3 to view the simulation pipeline diagram and method 4 to view the code hotspots for further analysis. \
	85	+ For more general cases such as MTE2 data transfer and actual scenarios: The shapes of the three input matrices are (128,128), (128,1), and (128,1), respectively, and the data type is float16. The current algorithm uses the two-pass method. Therefore, X is moved in for three times, and W and B are moved in for one time. The total amount of data to be transferred can be calculated accordingly. The theoretical value calculated based on the method described in the [Theoretical Parameters](#theoretical-parameters) section is sizeof(float16) * (128 * 128 * 3 + 128 + 128)/1.8 TB/s ≈ 0.1991 μs (calculated based on 1 TB = 10<sup>12</sup> Byte), which is greatly different from the actual performance data aiv_mte2_time. Analysis shows the total input size is smaller than the Unified Buffer (UB) capacity (192 KB for the A2 model). Therefore, if the MTE2 time is excessive, the basic block obtained through tiling computation may be too small, triggering redundant transfer instructions. In this case, pipeline optimization and tiling optimization are required, you can refer to method 3 to view the simulation pipeline diagram and analyze each pipeline for further analysis.
	86	+
	87	+- Method 2: Use board profiling to analyze the tiling.
	88	+The AI processor used in the previous example has 48 vector cores. The _layer_norm_fwd_fused operator is a pure vector operator. However, in some scenarios, too many blocks (Block Dim > 48) are delivered, causing excessive host scheduling overhead. In this case, the next step is to optimize the tiling.
	89	+
	90	+- Method 3: Use the simulation pipeline diagram to analyze the pipeline.
	91	+![analyse_data_waveform](../figures/performance_analysis_analyse_data_waveform.png) \
	92	+ The preceding figure shows the data obtained from an AI processor simulator. It can be seen that the SCALAR and FLOWCTRL instructions of the Vector core are saturated. You can analyze the operator logic to check whether there are too many scalar computations and unsupported vectorization operations. The next step is to optimize scalar computation. On the other hand, the related pipelines (such as MTE2 and VECTOR of veccore0) of the Vector core are regularly interrupted, that is, there are a large number of blank segments without operations. You can analyze the operator logic to check whether the stream interruption is caused by small basic block splitting. The main optimization direction is pipeline optimization. In addition, the vector pipeline utilization is further improved using tiling optimization and memory optimization.
	93	+
	94	+- Method 4: Analyze the code hotspot.
	95	+![analyse_data_code_mapping](../figures/performance_analysis_analyse_data_code_mapping.png) \
	96	+ The preceding figure shows the data obtained from an AI processor simulator. The load interface on the left corresponds to a group of assembly instructions on the right (only instructions related to code lines are displayed and sorted in descending order by cycle count). The high proportion of scalar instructions is inconsistent with the scenario where the MTE proportion should be high when load is used as the memory access interface. Therefore, the main optimization direction is scalar calculation.
	97	+
	98	+### Example: i64/i32 Comparison Failing to Vectorize on NPU, Leading to Scalar Fallback
	99	+
	100	+[Description] The i64/i32 comparison (cmp) cannot enable Vector on the NPU, causing them to degenerate into scalar computation and reducing efficiency. The i64/i32 cmp is converted to fp32 to accelerate vector operations by using vec_cast and vec_cmp.
	101	+[Note] When cmp is used within a mask in tl.load or tl.save, the compiler can typically auto-vectorize the operation. In this example, tl.where requires manual intervention to ensure vectorization.
	102	+
	103	+
	104	+```diff
	105	+@triton.jit
	106	+def npu_vector_cmp_kernel(
	107	+ X, # [Tensor] input tensor (row x col)
	108	+ Out, # [Tensor] output tensor (row x col)
	109	+ Mean, # [Vector] mean tensor (row, ) of X
	110	+ Rstd, # [Vector] std tensor (row, ) of X
	111	+ stride_x_row, # [Scalar] stride of row in X
	112	+ stride_out_row, # [Scalar] stride of row in Out, normally equals to stride_x_row
	113	+ M, # [Scalar] row number
	114	+ N, # [Scalar] col number
	115	+ eps, # [Scalar] epsilon to aviod division by zeros
	116	+ BLOCK_M: tl.constexpr,
	117	+ BLOCK_N: tl.constexpr
	118	+):
	119	+ """
	120	+ an example of layernorm to checkout Vector Cmp
	121	+ Out = ((X - E[X]) / sqrt(V[X] + eps)) on dim -1
	122	+
	123	+ just for easy case, we assume that:
	124	+ 1. BLOCK_N >= X.shape(-1), group_n = 0 only
	125	+ 2. BLOCK_M = 1, group_m = range(0, row, 1)
	126	+ """
	127	+ group_m = tl.program_id(0)
	128	+ group_n = tl.program_id(1)
	129	+ row = group_m
	130	+
	131	+ # calculate index & offset
	132	+ Mean = Mean + group_n * M
	133	+ Rstd = Rstd + group_n * M
	134	+ X = X + row * stride_x_row + group_n * N
	135	+ Out = Out + row * stride_out_row + group_n * N
	136	+
	137	+ cols = tl.arange(0, BLOCK_N) # cols is int64
	138	+ x = tl.load(X + cols, mask=cols < N, other=0.0).to(tl.float32)
	139	+
	140	+ # calculate mean & rstd
	141	+ mean = tl.sum(x, axis=0) / N
	142	+ tl.store(Mean + row, mean)
	143	+
	144	+- xbar = tl.where(cols < N, x - mean, 0.0) # N is a scalar value.
	145	+
	146	++ # change cols(i64) into cols_cmp(f32) to enable vector processing
	147	++ cols_cmp = cols.to(tl.float32)
	148	++ xbar = tl.where(cols_cmp < N, x - mean, 0.0)
	149	+
	150	+ var = tl.sum(xbar * xbar, axis=0) / N
	151	+ rstd = 1 / tl.sqrt(var + eps)
	152	+ tl.store(Rstd + row, rstd)
	153	+
	154	+ # calculate Out
	155	+ mask = cols < N
	156	+ out = (x - mean) * rstd
	157	+ tl.store(Out + cols, out, mask=mask)
	158	+```
	159	+Example Data comparison before and after optimization
	160	+![Figure 2 optimization2](../figures/optimization.png)
	161	+According to the data in the figure, the values of aiv_scalar_time (in μs) and aiv_scalar_ratio before and after optimization are greatly different, indicating that the performance is poor due to many scalar operations.
	162	+You can obtain visualize_data.bin by collecting the [operator simulation pipeline diagram](#operator-simulation-pipeline-diagram). Then, use MindStudio Insight to parse visualize_data.bin. It is found that xbar = tl.where(cols < N, x - mean, 0.0) contains many scalar operations, which can be reduced through the preceding optimization.

Adocs/en/environment_variable_reference.md+35-0

@@ -0,0 +1,35 @@
	1	+## Environment Variables
	2	+
	3	+The following table describes how to set environment variables.
	4	+
	5	+\| Category\| Environment Variable\| Default Value\| Function Description\| Setting Description\| Change Description\|
	6	+\|------\|----------\|--------\|----------\|----------\|----------\|
	7	+\| Debugging and logging\| TRITON_DEBUG \| 0 or not set\| Specifies whether to enable the debugging output function of Triton to print detailed debugging information during running. This is useful for troubleshooting problems in the compilation or execution phase. When this parameter is set to 1, Triton outputs more information about the compilation, kernel generation, and execution. Some implementations may support more fine-grained debugging levels (such as 2 and 3), depending on the Triton version and implementation.\| 0: The debugging is disabled.<br>1: The debugging is enabled.\| \|
	8	+\| Debugging and logging\| MLIR_ENABLE_DUMP \| 0 or not set\| Specifies whether to dump the intermediate representation (IR) of all kernels before each MLIR optimization. You can set `MLIR_ENABLE_DUMP` to `kernelName` to dump the IR of a specific kernel.\| 0: Do not dump.<br>1: Dump the IR of all kernels.<br>kernelName: Dump the IR of a specific kernel.\| The Triton cache may interfere with the dump. If `MLIR_ENABLE_DUMP=1` does not take effect, you can run `rm -r ~/.triton/cache/*` to clear the Triton cache.\|
	9	+\| Debugging and logging\| LLVM_IR_ENABLE_DUMP \| 0 or not set\| Specifies whether to dump the IR before each LLVM IR optimization.\| 0: Do not dump.<br>1: Dump IRs.\| \|
	10	+\| Debugging and logging\| TRITON_REPRODUCER_PATH \| Not set\| Generates the MLIR reproduction file before each MLIR compilation phase. If a phase fails, `<reproducer_path>` saves the MLIR status before the failure.\| `<reproducer_path>`: save path.\| \|
	11	+\| Debugging and logging\| TRITON_INTERPRET \| 0 or not set\| Specifies whether to use the Triton interpreter instead of the GPU for running and support inserting Python breakpoints in kernel function code.\| 0: Breakpoints are not supported.<br>1: Breakpoints are supported.\| \|
	12	+\| Debugging and logging\| TRITON_ENABLE_LLVM_DEBUG \| 0 or not set\| Specifies whether to pass the`-debug` parameter to LLVM and outputs a large amount of debugging information. If there is too much information, you can use `TRITON_LLVM_DEBUG_ONLY` to limit the output scope.\| 0: Pass.<br>1: Do not pass.\| Another method to reduce output interference is as follows: Set the running program by setting `LLVM_IR_ENABLE_DUMP` to `1`, extract the IR before the target LLVM optimization channel, and run the `opt` tool of the LLVM separately. In this case, you can add `-debug-only=foo` to the command line to limit the debugging range.\|
	13	+\| Debugging and logging\| TRITON_LLVM_DEBUG_ONLY \| Not set\| Equivalent to the `-debug-only` command line option of LLVM. This parameter can be used to limit the LLVM debugging output to a specific optimization channel or component name (defined by the `#define DEBUG_TYPE` macro in LLVM and Triton), thereby effectively reducing redundant debugging output. You can specify one or more comma-separated values, for example, `TRITON_LLVM_DEBUG_ONLY="tritongpu-remove-layout-conversions"` or `TRITON_LLVM_DEBUG_ONLY="tritongpu-remove-layout-conversions,regalloc"`.\| Comma-separated values: channel or component name\| \|
	14	+\| Debugging and logging\| USE_IR_LOC \| 0 or not set\| Specifies whether to include location information (such as file names and line numbers) in the generated IR. This information is helpful for debugging, but may increase the size of the generated IR. If this parameter is set to 1, the IR is re-parsed, and the location information is mapped to the line number of the IR file with a specific extension (not the line number of the Python source file). This enables a direct mapping from the IR to the LLVM IR/PTX. When used with the performance analysis tool, this parameter can be used to implement fine-grained performance analysis on IR instructions.\| 0: No location information is included.<br>1: The location information is included.\| \|
	15	+\| Debugging and logging\| TRITON_PRINT_AUTOTUNING \| 0 or not set\| After the automatic optimization is complete, the optimal configuration and total time of each kernel are output.\| 0: Do not output.<br>1: Output.\| \|
	16	+\| Debugging and logging\| MLIR_ENABLE_REMARK \| 0 or not set\| Specifies whether to enable the output of remarks during MLIR compilation, including performance warnings in remarks.\| 0: Disabled.<br>1: Enabled.\| \|
	17	+\| Debugging and logging\| TRITON_KERNEL_DUMP \| 0 or not set\| Specifies whether to enable the dump function of the Triton kernel. When this function is enabled, Triton saves the generated kernel code (IR and final PTX in each compilation phase) to the specified directory.\| 0: Disabled.<br>1: Enabled.\| \|
	18	+\| Debugging and logging\| TRITON_DUMP_DIR \| Current working directory or not set\| Specifies the directory for storing the Triton kernel dump file, which is the directory for saving the IR and PTX when `TRITON_KERNEL_DUMP` is set to `1`.\| "path": save path.\| \|
	19	+\| Debugging and logging\| TRITON_DEVICE_PRINT \| 0 or not set\| If this parameter is set to `1` or `true` (`TRUE` is converted to `true`), the function of `tl.device_print` is enabled. Note: This function uses the GM buffer (the pointer of which is passed to the kernel).\| 0: Disabled.<br>1: The functionality of `tl.device_print` is enabled.\| The maximum size of the GM buffer for each thread is 16 KB. If the buffer size exceeds 16 KB, the excess content will be discarded. The value is fixed currently and will be adjusted through an environment variable.\|
	20	+\| Compilation control\| TRITON_ALWAYS_COMPILE \| 0 or not set\| Specifies whether Triton forcibly recompiles the kernel each time it runs, instead of using the existing cached version. By default, Triton caches the compiled kernels (based on parameters and configurations) to improve performance. If this parameter is set to 1, Triton ignores the cache and recompiles the kernel each time it runs, which is useful for debugging or testing new compiler features.\| 0: Disabled.<br>1: All kernels are recompiled during each running.\| \|
	21	+\| Compilation control\| DISABLE_LLVM_OPT \| 0 or not set\| If this parameter is set to 1, the optimization steps (LLVM optimization of make_llir and make_ptx) during LLVM compilation can be disabled. If this parameter is set to a character string, the LLVM optimization flags to be disabled are parsed. For example, if `DISABLE_LLVM_OPT` is set to `"disable-lsr"`, the loop strength optimization is disabled (this optimization may cause a performance fluctuation of up to 10% in some kernels with register pressure).\| 0: The LLVM optimization is enabled.<br>1: The optimization steps (LLVM optimization of make_llir and make_ptx) during LLVM compilation are disabled.\| \|
	22	+\| Compilation control\| MLIR_ENABLE_TIMING \| 0 or not set\| Specifies whether to enable the time statistics function during MLIR compilation.\| 0: Disabled.<br>1: Enabled.\| \|
	23	+\| Compilation control\| LLVM_ENABLE_TIMING \| 0 or not set\| Specifies whether to enable the time statistics function during LLVM compilation.\| 0: Disabled.<br>1: Enabled.\| \|
	24	+\| Compilation control\| TRITON_DEFAULT_FP_FUSION \| 1 (enabled)\| Specifies whether to enable the floating-point operation fusion optimization by default. The default floating-point operation fusion behavior (for example, mul+add->fma) is overwritten.\| 0: Disabled.<br>1: Enabled.\| \|
	25	+\| Compilation control\| TRITON_KERNEL_OVERRIDE \| 0 or not set\| Specifies whether to enable the Triton kernel override function. You can use the user-specified external file (such as IR/PTX) to override the default generated kernel code at the beginning of each compilation phase.\| 0: Disabled.<br>1: Enabled.\| \|
	26	+\| Compilation control\| TRITON_OVERRIDE_DIR \| Current working directory or not set\| Specifies the directory for searching the Triton kernel override file. Directory for loading the IR/PTX file when `TRITON_KERNEL_OVERRIDE` is set to `1`.\| "path": save path.\| \|
	27	+\| Compilation control\| TRITON_ASCEND_COMPILE_SPEED_OPT \| 0 or not set\| Specifies whether the JIT compiler skips the subsequent compilation phase after detecting that the kernel compilation fails. Set the parameter to `1` to skip the attempt. (The default value `0` indicates that the attempt is continued.)\| 0: Continue the attempt.<br>1: Skip.\| \|
	28	+\| Compilation control\| TRITON_COMPILE_ONLY \| 0 or not set\| Specifies whether to perform only compilation without execution. This parameter is used when remote_launch is used.\| 0: The optimization is disabled.<br>1: Enabled.\| \|
	29	+\| Compilation control\| TRITON_DISABLE_FFTS \| 0 or not set\| Specifies whether to enable FFTS.\| 0: The optimization is disabled.<br>1: Enabled.\| \|
	30	+\| Running and scheduling\| TRITON_ALL_BLOCKS_PARALLEL \| 0 or not set\| Specifies whether to enable the automatic optimization of the number of logical cores based on the number of physical cores. This parameter can be enabled only when logical cores can execute in parallel. When the number of logical cores is greater than the number of physical cores, enabling this parameter will instruct the compiler to automatically adjust the number of logical cores to match the number of physical cores, thereby reducing scheduling overhead. After this parameter is enabled, the value of grid can be greater than 65535. Limitation: This option can be enabled only when the logic of the Triton kernel is insensitive to the execution sequence. Otherwise, a deadlock may occur.\| 0: The optimization is disabled.<br>1: Enabled.\| \|
	31	+\| Running and scheduling\| TRITON_ENABLE_TASKQUEUE \| 0 or not set\| Specifies whether to enable task_queue.\| 0: The optimization is disabled.<br>1: Enabled.\| \|
	32	+\| Running and scheduling\| TRITON_ENABLE_SANITIZER \| 0 or not set\| Specifies whether to enable SANITIZER.\| 0: The optimization is disabled.<br>1: Enabled.\| \|
	33	+\| Running and scheduling\| ENABLE_PRINT_UB_BITS \| 0 or not set\| After this parameter is enabled, the current UB usage can be obtained for the inductor.\| 0: The optimization is disabled.<br>1: Enabled.\| \|
	34	+\| Others\| TRITON_BENCH_METHOD \| Not set\| When the Ascend NPU is used, change `do_bench` in `testing.py` to `do_bench_npu`. (This parameter is used when `INDUCTOR_ASCEND_AGGRESSIVE_AUTOTUNE` is set to `1`.) If this parameter is set to `default`, the original `do_bench` function is still called even if the NPU is available.\| "npu": Switch to `do_bench_npu`.\| \|
	35	+\| Others\| TRITON_REMOTE_RUN_CONFIG_PATH \| path \| Specifies the configuration path for remote running.\| Specify the path directly.\| \|

Adocs/en/examples/01_vector_add_example.md+90-0

@@ -0,0 +1,90 @@
	1	+# Vector Addition
	2	+
	3	+In this section, you will use Triton to write a simple vector addition program.
	4	+In this process, you will learn:
	5	+
	6	+- The basic programming model of Triton.
	7	+- The `triton.jit` decorator used to define Triton kernels.
	8	+
	9	+Compute kernel:
	10	+
	11	+```bash
	12	+import torch
	13	+import torch_npu
	14	+
	15	+import triton
	16	+import triton.language as tl
	17	+
	18	+
	19	+@triton.jit
	20	+def add_kernel(x_ptr, # Pointer to the first input vector.
	21	+ y_ptr, # Pointer to the second input vector.
	22	+ output_ptr, # Pointer to the output vector.
	23	+ n_elements, # Size of the vector.
	24	+ BLOCK_SIZE: tl.constexpr, # Number of elements that should be processed by each program.
	25	+ # Note: `constexpr` will mark the variable as a constant.
	26	+ ):
	27	+ # Different data is processed by different "processes", so you need to allocate:
	28	+ pid = tl.program_id(axis=0) # A 1D launch grid is used, so the axis is 0.
	29	+ # This program will process inputs that are offset from the initial data.
	30	+ # For example, if there is a vector of length 256 and block size 64, the program will access the elements [0:64, 64:128, 128:192, 192:256] respectively.
	31	+ # Note that offsets is a list of pointers:
	32	+ block_start = pid * BLOCK_SIZE
	33	+ offsets = block_start + tl.arange(0, BLOCK_SIZE)
	34	+ # Create a mask to prevent memory operations from out-of-bounds accesses.
	35	+ mask = offsets < n_elements
	36	+ # Load x and y from DRAM, and mask out any extra elements if the input is not an integer multiple of the block size.
	37	+ x = tl.load(x_ptr + offsets, mask=mask)
	38	+ y = tl.load(y_ptr + offsets, mask=mask)
	39	+ output = x + y
	40	+ # Write x + y back to DRAM.
	41	+ tl.store(output_ptr + offsets, output, mask=mask)
	42	+```
	43	+
	44	+Create a helper function to:
	45	+
	46	+- Generate the z tensor;
	47	+- Enqueue the above kernel with the appropriate grid/block sizes.
	48	+
	49	+```Python
	50	+def add(x: torch.Tensor, y: torch.Tensor):
	51	+ # The output needs to be pre-allocated.
	52	+ output = torch.empty_like(x)
	53	+ n_elements = output.numel()
	54	+ # The launch grid indicates the number of kernel instances that run in parallel.
	55	+ # It can be Tuple[int] or Callable(metaparameters) -> Tuple[int].
	56	+ # In this case, a 1D grid is used, where the size is the number of blocks:
	57	+ grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']), )
	58	+ # NOTE:
	59	+ # - Each torch.tensor object is implicitly converted into a pointer to its first element.
	60	+ # - The `triton.jit` function can be indexed with a launch grid to obtain a callable GPU kernel.
	61	+ # - Pass meta-parameters as keywords.
	62	+ add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=1024)
	63	+ # Returns the handle to z.
	64	+ return output
	65	+```
	66	+
	67	+Use the above function to compute the element-wise sum of two `torch.tensor` objects and test its correctness:
	68	+
	69	+```Python
	70	+torch.manual_seed(0)
	71	+size = 98432
	72	+x = torch.rand(size, device='npu')
	73	+y = torch.rand(size, device='npu')
	74	+output_torch = x + y
	75	+output_triton = add(x, y)
	76	+print(output_torch)
	77	+print(output_triton)
	78	+print(f'The maximum difference between torch and triton is '
	79	+ f'{torch.max(torch.abs(output_torch - output_triton))}')
	80	+```
	81	+
	82	+Output:
	83	+
	84	+```bash
	85	+tensor([0.8329, 1.0024, 1.3639, ..., 1.0796, 1.0406, 1.5811], device='npu:0')
	86	+tensor([0.8329, 1.0024, 1.3639, ..., 1.0796, 1.0406, 1.5811], device='npu:0')
	87	+The maximum difference between torch and triton is 0.0
	88	+```
	89	+
	90	+"The maximum difference between torch and triton is 0.0" indicates that the output results of Triton and PyTorch are the same.

Adocs/en/examples/02_fused_softmax_example.md+151-0

@@ -0,0 +1,151 @@
	1	+# Fused Softmax
	2	+
	3	+In this section, you will use Triton to write a program of the fused softmax operation.
	4	+In this process, you will learn:
	5	+
	6	+- The advantages of kernel fusion for bandwidth-bound operations.
	7	+- Reduction operations in Triton.
	8	+
	9	+## Using Native PyTorch to Perform Softmax Operation on X Row by Row
	10	+
	11	+```Python
	12	+import torch
	13	+import torch_npu
	14	+
	15	+import triton
	16	+import triton.language as tl
	17	+
	18	+def naive_softmax(x):
	19	+ """
	20	+ Subtract the maximum element to avoid overflow. Softmax is invariant to this offset.
	21	+ """
	22	+ # Read MN elements; write M elements.
	23	+ x_max = x.max(dim=1)[0]
	24	+ # Read MN + M elements; write MN elements.
	25	+ z = x - x_max[:, None]
	26	+ # Read MN elements; write MN elements.
	27	+ numerator = torch.exp(z)
	28	+ # Read MN elements; write M elements.
	29	+ denominator = numerator.sum(dim=1)
	30	+ # Read MN + M elements; write MN elements.
	31	+ ret = numerator / denominator[:, None]
	32	+ # Total: Read 5 × MN + 2 × M elements; write 3 × MN + 2 × M elements.
	33	+ return ret
	34	+```
	35	+
	36	+Purpose of kernel fusion
	37	+
	38	+When implemented naively in PyTorch, computing `y = naive_softmax(x)` requires reading 5 × MN + 2 × M elements from DRAM and writing back 3 MN + 2 M elements. Obviously, this is very inefficient. A more efficient solution is to use a custom "fused" kernel that reads `x` only once and completes all necessary computations on the chip.
	39	+Doing so requires reading and writing back only 2 × MN bytes. Therefore, the theoretical speedup ratio is about 4 times, that is, 8 × MN + 4 × M)/2 × MN.
	40	+
	41	+`torch.jit.script` is designed to automatically perform this kind of "kernel fusion", but it is still far from ideal.
	42	+
	43	+## Compute Kernel
	44	+
	45	+The softmax kernel works as follows: Each compute unit (program) loads a group of data rows of the input matrix X stridden by number of programs, normalizes it, and writes back the result to the output matrix Y.
	46	+Note: A significant limitation of Triton is that each block must have a power-of-two number of elements. Therefore, to handle any possible input shapes, internally "pad" each row and ensure the correctness of memory operations.
	47	+
	48	+```Python
	49	+@triton.jit
	50	+def softmax_kernel(output_ptr, input_ptr, input_row_stride, output_row_stride, n_rows, n_cols, BLOCK_SIZE: tl.constexpr):
	51	+ # Program start row
	52	+ row_start = tl.program_id(0)
	53	+ row_step = tl.num_programs(0)
	54	+ for row_idx in tl.range(row_start, n_rows, row_step):
	55	+ # The stride indicates the required increment of the pointer to advance one row.
	56	+ row_start_ptr = input_ptr + row_idx * input_row_stride
	57	+ # The block size is the next power of two greater than n_cols, so we can fit
	58	+ # rows in a single block.
	59	+ col_offsets = tl.arange(0, BLOCK_SIZE)
	60	+ input_ptrs = row_start_ptr + col_offsets
	61	+ # Load the row into SRAM using a mask, because BLOCK_SIZE may be greater than n_cols.
	62	+ mask = col_offsets < n_cols
	63	+ row = tl.load(input_ptrs, mask=mask, other=-float('inf'))
	64	+ # Subtract the maximum value for numerical stability.
	65	+ row_minus_max = row - tl.max(row, axis=0)
	66	+ # Note that exponentiation in Triton is fast but approximate.
	67	+ numerator = tl.exp(row_minus_max)
	68	+ denominator = tl.sum(numerator, axis=0)
	69	+ softmax_output = numerator / denominator
	70	+ # Write the output back to DRAM.
	71	+ output_row_start_ptr = output_ptr + row_idx * output_row_stride
	72	+ output_ptrs = output_row_start_ptr + col_offsets
	73	+ tl.store(output_ptrs, softmax_output, mask=mask)
	74	+```
	75	+
	76	+Create a helper function. This function can add the kernel function and its meta-parameters to the execution queue to process any given input tensor.
	77	+
	78	+```Python
	79	+target = triton.runtime.driver.active.get_current_target()
	80	+kernels = {}
	81	+
	82	+def softmax(x, stream):
	83	+ n_rows, n_cols = x.shape
	84	+
	85	+ # The block size for each loop iteration is the smallest power of two greater than or equal to the number of columns in `x`.
	86	+ BLOCK_SIZE = triton.next_power_of_2(n_cols)
	87	+ # Allocate output space.
	88	+ y = torch.empty_like(x)
	89	+
	90	+ # Precompile the kernel to obtain the register usage and compute the thread occupancy.
	91	+ kernel, num_programs = kernels.get(BLOCK_SIZE, (None, 0))
	92	+ if kernel is None:
	93	+ num_programs = 32
	94	+ kernel = softmax_kernel
	95	+ kernels[BLOCK_SIZE] = (kernel, num_programs)
	96	+
	97	+ num_programs = min(num_programs, n_rows)
	98	+
	99	+ kernel[(num_programs, 1, 1)](
	100	+ y,
	101	+ x,
	102	+ x.stride(0),
	103	+ y.stride(0),
	104	+ n_rows,
	105	+ n_cols,
	106	+ BLOCK_SIZE
	107	+ )
	108	+ return y
	109	+```
	110	+
	111	+## Unit Test
	112	+
	113	+The processed kernel needs to be tested on a matrix with irregular numbers of rows and columns. This can verify that the padding mechanism works.
	114	+
	115	+```Python
	116	+device = torch.npu.current_device()
	117	+stream = torch.npu.current_stream(device).npu_stream
	118	+torch.manual_seed(0)
	119	+x = torch.randn(1823, 781, device='npu')
	120	+y_triton = softmax(x, stream)
	121	+y_torch = torch.softmax(x, axis=1)
	122	+assert torch.allclose(y_triton, y_torch), (y_triton, y_torch)
	123	+print(y_triton)
	124	+print(y_torch)
	125	+print(f'The maximum difference between torch and triton is '
	126	+ f'{torch.max(torch.abs(y_triton-y_torch))}')
	127	+```
	128	+
	129	+Output:
	130	+
	131	+```bash
	132	+tensor([[0.0002, 0.0017, 0.0009, ..., 0.0009, 0.0013, 0.0073],
	133	+ [0.0001, 0.0004, 0.0006, ..., 0.0006, 0.0004, 0.0003],
	134	+ [0.0007, 0.0002, 0.0006, ..., 0.0011, 0.0004, 0.0039],
	135	+ ...,
	136	+ [0.0021, 0.0002, 0.0015, ..., 0.0012, 0.0014, 0.0022],
	137	+ [0.0003, 0.0002, 0.0007, ..., 0.0005, 0.0006, 0.0007],
	138	+ [0.0034, 0.0014, 0.0005, ..., 0.0007, 0.0016, 0.0028]],
	139	+ device='npu:0')
	140	+tensor([[0.0002, 0.0017, 0.0009, ..., 0.0009, 0.0013, 0.0073],
	141	+ [0.0001, 0.0004, 0.0006, ..., 0.0006, 0.0004, 0.0003],
	142	+ [0.0007, 0.0002, 0.0006, ..., 0.0011, 0.0004, 0.0039],
	143	+ ...,
	144	+ [0.0021, 0.0002, 0.0015, ..., 0.0012, 0.0014, 0.0022],
	145	+ [0.0003, 0.0002, 0.0007, ..., 0.0005, 0.0006, 0.0007],
	146	+ [0.0034, 0.0014, 0.0005, ..., 0.0007, 0.0016, 0.0028]],
	147	+ device='npu:0')
	148	+The maximum difference between torch and triton is 1.4901161193847656e-08
	149	+```
	150	+
	151	+"The maximum difference between torch and triton is 1.4901161193847656e-08" indicates that the output results of Triton and PyTorch are very close and cannot be visually distinguished.

Adocs/en/examples/03_layer_norm_example.md+173-0

@@ -0,0 +1,173 @@
	1	+# Layer Normalization
	2	+
	3	+In this section, you will use Triton to write a high-performance layer normalization kernel that runs faster than the PyTorch implementation.
	4	+
	5	+## Compute Kernel
	6	+
	7	+```Python
	8	+import pytest
	9	+import torch
	10	+import triton
	11	+import triton.language as tl
	12	+import torch_npu
	13	+
	14	+@triton.jit
	15	+def _layer_norm_fwd_fused(
	16	+ X, # Pointer to the input
	17	+ Y, # Pointer to the output
	18	+ W, # Pointer to the weights
	19	+ B, # Pointer to the biases
	20	+ Mean, # Pointer to the mean
	21	+ Rstd, # Pointer to the 1/std
	22	+ stride, # Number of elements to be added when the pointer moves by one row
	23	+ N, # Number of columns in X
	24	+ eps, # Epsilon used to avoid division by zero
	25	+ BLOCK_SIZE: tl.constexpr,
	26	+):
	27	+ # Map the program ID to the corresponding rows of X and Y for computation.
	28	+ row = tl.program_id(0)
	29	+ Y += row * stride
	30	+ X += row * stride
	31	+ # Calculate the mean.
	32	+ mean = 0
	33	+ _mean = tl.zeros([BLOCK_SIZE], dtype=tl.float32)
	34	+ for off in range(0, N, BLOCK_SIZE):
	35	+ cols = off + tl.arange(0, BLOCK_SIZE)
	36	+ a = tl.load(X + cols, mask=cols < N, other=0.).to(tl.float32)
	37	+ _mean += a
	38	+ mean = tl.sum(_mean, axis=0) / N
	39	+ # Calculate the variance.
	40	+ _var = tl.zeros([BLOCK_SIZE], dtype=tl.float32)
	41	+ for off in range(0, N, BLOCK_SIZE):
	42	+ cols = off + tl.arange(0, BLOCK_SIZE)
	43	+ x = tl.load(X + cols, mask=cols < N, other=0.).to(tl.float32)
	44	+ x = tl.where(cols < N, x - mean, 0.)
	45	+ _var += x * x
	46	+ var = tl.sum(_var, axis=0) / N
	47	+ rstd = 1 / tl.sqrt(var + eps)
	48	+ # Write mean/rstd.
	49	+ tl.store(Mean + row, mean)
	50	+ tl.store(Rstd + row, rstd)
	51	+ # Normalize and apply linear transformation.
	52	+ for off in range(0, N, BLOCK_SIZE):
	53	+ cols = off + tl.arange(0, BLOCK_SIZE)
	54	+ mask = cols < N
	55	+ w = tl.load(W + cols, mask=mask)
	56	+ b = tl.load(B + cols, mask=mask)
	57	+ x = tl.load(X + cols, mask=mask, other=0.).to(tl.float32)
	58	+ x_hat = (x - mean) * rstd
	59	+ y = x_hat * w + b
	60	+ # Write the output.
	61	+ tl.store(Y + cols, y, mask=mask)
	62	+```
	63	+
	64	+LayerNorm Implementation Defined by Using Triton
	65	+
	66	+```Python
	67	+@torch.inference_mode()
	68	+def layer_norm(x, weight, bias, eps=1e-5):
	69	+ # Allocate the output tensor with the same shape and data type as the input.
	70	+ y = torch.empty_like(x)
	71	+
	72	+ # Flatten the input x into a two-dimensional shape [-1, feature_dim] for processing the last dimension.
	73	+ x_arg = x.reshape(-1, x.shape[-1])
	74	+ M, N = x_arg.shape
	75	+
	76	+ mean = torch.empty((M, ), dtype=torch.float32, device=x.device)
	77	+ rstd = torch.empty((M, ), dtype=torch.float32, device=x.device)
	78	+
	79	+ BLOCK_SIZE = 1024
	80	+
	81	+ # enqueue kernel
	82	+ kernel = _layer_norm_fwd_fused[(M,)](# M indicates the number of blocks, and launch grid=(M,)
	83	+ x_arg, y, weight, bias, mean, rstd, # Input, output, and intermediate variables
	84	+ x_arg.stride(0), N, eps,
	85	+ BLOCK_SIZE=BLOCK_SIZE)
	86	+ # Return the normalized output.
	87	+ return y
	88	+
	89	+# Call layer normalization during forward pass.
	90	+def _layer_norm(M, N, dtype, eps=1e-5, device='npu'):
	91	+ # Construct data.
	92	+ x_shape = (M, N)
	93	+ w_shape = (x_shape[-1], )
	94	+ weight = torch.rand(w_shape, dtype=dtype, device=device, requires_grad=True)
	95	+ bias = torch.rand(w_shape, dtype=dtype, device=device, requires_grad=True)
	96	+ x = -2.3 + 0.5 * torch.randn(x_shape, dtype=dtype, device=device)
	97	+ dy = .1 * torch.randn_like(x)
	98	+ x.requires_grad_(True)
	99	+ # Forward pass
	100	+ y_tri = layer_norm(x, weight, bias, eps)
	101	+ y_ref = torch.nn.functional.layer_norm(x, weight, bias, eps).to(dtype)
	102	+ # Determine whether the results are approximate.
	103	+ assert torch.allclose(y_tri, y_ref, atol=1e-2, rtol=0)
	104	+ print(f"y_tri: {y_tri}")
	105	+ print(f"y_ref: {y_ref}")
	106	+ print(f"Layer Normalization {M},{N} {dtype} PASSED!")
	107	+
	108	+# Perform the test.
	109	+if __name__ == '__main__':
	110	+ _layer_norm(128, 128, torch.float16)
	111	+ _layer_norm(128, 128, torch.bfloat16)
	112	+ _layer_norm(128, 128, torch.float32)
	113	+```
	114	+
	115	+Result
	116	+
	117	+```bash
	118	+y_tri: tensor([[ 0.2512, 0.0647, 0.8389, ..., 2.3652, 1.5039, 1.1904],
	119	+ [ 1.0908, 1.5391, 0.2269, ..., 1.6846, 1.0996, 0.9614],
	120	+ [-0.2974, 0.5918, 0.3225, ..., 2.2891, -0.8418, 0.6885],
	121	+ ...,
	122	+ [ 0.5225, -0.0068, 0.4968, ..., -1.1221, 1.7422, 0.6143],
	123	+ [ 0.4463, 1.2441, 0.2224, ..., 2.2969, -0.3311, 0.6177],
	124	+ [-0.0113, 0.8423, 0.3696, ..., 1.3838, 1.2471, 0.8750]],
	125	+ device='npu:0', dtype=torch.float16)
	126	+y_ref: tensor([[ 0.2512, 0.0647, 0.8389, ..., 2.3652, 1.5039, 1.1904],
	127	+ [ 1.0908, 1.5391, 0.2269, ..., 1.6846, 1.0996, 0.9614],
	128	+ [-0.2974, 0.5918, 0.3225, ..., 2.2891, -0.8418, 0.6885],
	129	+ ...,
	130	+ [ 0.5225, -0.0068, 0.4968, ..., -1.1221, 1.7422, 0.6143],
	131	+ [ 0.4463, 1.2441, 0.2224, ..., 2.2969, -0.3311, 0.6177],
	132	+ [-0.0113, 0.8423, 0.3696, ..., 1.3838, 1.2471, 0.8750]],
	133	+ device='npu:0', dtype=torch.float16, grad_fn=<NativeLayerNormBackward0>)
	134	+Layer Normalization 128,128 torch.float16 PASSED!
	135	+y_tri: tensor([[-0.4180, 0.9648, 0.8633, ..., 0.7656, 0.8438, 0.3633],
	136	+ [ 0.4453, 0.5352, 0.9102, ..., 1.1875, -0.0562, 0.5391],
	137	+ [ 1.3125, 0.9961, 0.9219, ..., 0.9688, 0.0025, 0.5156],
	138	+ ...,
	139	+ [-0.1426, 0.6289, 0.9609, ..., 0.9648, -0.1260, -0.1270],
	140	+ [ 1.1641, 0.6680, 0.8281, ..., 0.9258, 0.9062, 0.1768],
	141	+ [-0.2129, 0.7109, 0.9141, ..., 0.7891, -0.0767, 0.5156]],
	142	+ device='npu:0', dtype=torch.bfloat16)
	143	+y_ref: tensor([[-0.4180, 0.9648, 0.8633, ..., 0.7656, 0.8438, 0.3633],
	144	+ [ 0.4453, 0.5352, 0.9102, ..., 1.1875, -0.0562, 0.5391],
	145	+ [ 1.3125, 0.9961, 0.9219, ..., 0.9688, 0.0025, 0.5156],
	146	+ ...,
	147	+ [-0.1426, 0.6289, 0.9609, ..., 0.9648, -0.1260, -0.1270],
	148	+ [ 1.1641, 0.6680, 0.8281, ..., 0.9258, 0.9062, 0.1768],
	149	+ [-0.2129, 0.7109, 0.9141, ..., 0.7891, -0.0767, 0.5156]],
	150	+ device='npu:0', dtype=torch.bfloat16, grad_fn=<NativeLayerNormBackward0>)
	151	+Layer Normalization 128,128 torch.bfloat16 PASSED!
	152	+y_tri: tensor([[-0.2980, 0.2922, 0.6481, ..., 0.9786, 0.7304, 0.8982],
	153	+ [ 1.5911, 0.0474, 0.6518, ..., 0.8013, 0.2435, 1.3748],
	154	+ [ 1.3024, 0.6265, 0.6473, ..., 0.8423, 0.0984, -1.1839],
	155	+ ...,
	156	+ [-0.2195, 0.1359, 0.6461, ..., 0.8319, 1.0899, 1.5015],
	157	+ [ 0.6371, 0.3687, 0.6530, ..., 0.9359, 0.0818, 0.6499],
	158	+ [ 0.1178, 0.3639, 0.6475, ..., 0.7221, 0.4622, 1.4510]],
	159	+ device='npu:0')
	160	+y_ref: tensor([[-0.2980, 0.2922, 0.6481, ..., 0.9786, 0.7304, 0.8982],
	161	+ [ 1.5911, 0.0474, 0.6518, ..., 0.8013, 0.2435, 1.3748],
	162	+ [ 1.3024, 0.6265, 0.6473, ..., 0.8423, 0.0984, -1.1839],
	163	+ ...,
	164	+ [-0.2195, 0.1359, 0.6461, ..., 0.8319, 1.0899, 1.5015],
	165	+ [ 0.6371, 0.3687, 0.6530, ..., 0.9359, 0.0818, 0.6499],
	166	+ [ 0.1178, 0.3639, 0.6475, ..., 0.7221, 0.4622, 1.4510]],
	167	+ device='npu:0', grad_fn=<NativeLayerNormBackward0>)
	168	+Layer Normalization 128,128 torch.float32 PASSED!
	169	+```
	170	+
	171	+"Layer Normalization 128,128 torch.float16 PASSED!", \
	172	+"Layer Normalization 128,128 torch.bfloat16 PASSED!", \
	173	+The result "Layer Normalization 128,128 torch.float32 PASSED!" indicates that the output of float16, bfloat16, and float32 data types on Triton is the same as that on PyTorch.

Adocs/en/examples/04_fused_attention_example.md+359-0

@@ -0,0 +1,359 @@
	1	+# Fused Attention
	2	+
	3	+This section implements a fused attention forward pass kernel of the Flash Attention v2 style based on Triton, which is applicable to the Ascend NPU platform. The implementation supports:
	4	+- Causal and non-causal attention
	5	+- Tiling for processing long sequences
	6	+- Max-shifted softmax for numerical stability optimization
	7	+
	8	+The overall structure contains two core Triton kernels:
	9	+1. `_attn_fwd_inner`: performs attention computation between a single query block and key/value blocks (causal masks are processed in phases).
	10	+2. `_attn_fwd`: schedules all query blocks and manages the block pointer, accumulator, and normalization.
	11	+
	12	+The `attention` function is encapsulated as a callable function using PyTorch `autograd.Function` and is verified for precision alignment with `torch_npu.npu_fusion_attention`.
	13	+
	14	+```Python
	15	+import pytest
	16	+import torch
	17	+import torch_npu
	18	+import triton
	19	+import triton.language as tl
	20	+
	21	+
	22	+DEVICE = "npu"
	23	+
	24	+
	25	+@triton.jit
	26	+def _attn_fwd_inner(acc_ptr, l_i, m_i, q, # Accumulator, local l, local m, query vector
	27	+ K_block_ptr, V_block_ptr, # Key and value block pointers for current stage
	28	+ start_m, qk_scale, # Starting position of current query block, qk scale factor
	29	+ BLOCK_M: tl.constexpr, HEAD_DIM: tl.constexpr, BLOCK_N: tl.constexpr, # Block size constants
	30	+ STAGE: tl.constexpr, offs_m: tl.constexpr, offs_n: tl.constexpr, # Current stage flag, m and n offset indices
	31	+ N_CTX: tl.constexpr, fp8_v: tl.constexpr): # Total context length, whether to enable FP8 for value precision
	32	+ # Set the processing range [lo, hi) for the current stage (in column block units)
	33	+ # causal = true
	34	+ # stage = 1
	35	+ # Causal attention, as the name implies, restricts the flow of information during computation,
	36	+ # only allowing the model to see the current and previous positions.
	37	+ # In other words, the output at the current position can only depend on the input at or before this position,
	38	+ # and cannot access information from future positions.
	39	+ # Causal attention ensures sequential order and prevents "leakage of future information."
	40	+ # But the following logic will also be triggered
	41	+ if STAGE == 1:
	42	+ # Stage 1: process all tokens before the query block
	43	+ tl.static_assert(BLOCK_M >= BLOCK_N)
	44	+ lo, hi = 0, start_m * BLOCK_M
	45	+ elif STAGE == 2:
	46	+ # Stage 2: process the current query block
	47	+ tl.static_assert(BLOCK_M >= BLOCK_N)
	48	+ lo, hi = start_m * BLOCK_M, (start_m + 1) * BLOCK_M
	49	+ lo = tl.multiple_of(lo, BLOCK_M) # Align starting position
	50	+ # causal = False (no need for masking)
	51	+ else:
	52	+ lo, hi = 0, N_CTX # Process the entire context
	53	+
	54	+ # Adjust K and V block pointers to the starting position `lo`
	55	+ K_block_ptr = tl.advance(K_block_ptr, (lo, 0)) # K is [HEAD_DIM, N_CTX], shift along the second dim by lo
	56	+ V_block_ptr = tl.advance(V_block_ptr, (lo, 0)) # V is [N_CTX, HEAD_DIM], shift along the first dim by lo
	57	+
	58	+ # Index mapping for the accumulator , used for slicing when HEAD_DIM >= 256
	59	+ row = tl.arange(0, BLOCK_M)[:, None]
	60	+ col_head_dim = tl.arange(0, HEAD_DIM)[None, :]
	61	+ block2d_acc = row * HEAD_DIM + col_head_dim
	62	+
	63	+ # Iterate over all k, v blocks in the current stage and accumulate the output
	64	+ for start_n in range(lo, hi, BLOCK_N): # Process BLOCK_N columns at a time
	65	+ start_n = tl.multiple_of(start_n, BLOCK_N) # Align column start position
	66	+ # -- Compute qk ----
	67	+ k = tl.load(K_block_ptr)
	68	+ # Modify K
	69	+ trans_k = tl.trans(k)
	70	+ qk = tl.dot(q, trans_k)
	71	+ # Apply causal mask for STAGE 2
	72	+ if STAGE == 2:
	73	+ mask = offs_m[:, None] >= (start_n + offs_n[None, :]) # Construct upper triangular mask
	74	+ qk = qk * qk_scale + tl.where(mask, 0, -1.0e6) # Set invalid positions to -∞
	75	+ m_ij = tl.maximum(m_i, tl.max(qk, 1)) # Update m_ij = max(m_i, max(qk))
	76	+ qk -= m_ij[:, None] # Subtract max for softmax stability
	77	+ else:
	78	+ qk = qk * qk_scale
	79	+ m_ij = tl.maximum(m_i, tl.max(qk, 1)) # Scaled max
	80	+ qk = qk - m_ij[:, None] # Stabilize
	81	+
	82	+ # Softmax weights p = exp(qk)
	83	+ p = tl.math.exp(qk)
	84	+
	85	+ # Convert softmax weight type depending on FP8 usage
	86	+ if fp8_v:
	87	+ p_cast = p.to(tl.float8e5) # Convert to FP8 format (save memory)
	88	+ else:
	89	+ p_cast = p.to(k.dtype)
	90	+
	91	+ v = tl.load(V_block_ptr) # Load corresponding V block
	92	+ pv = tl.dot(p_cast, v)
	93	+ l_ij = tl.sum(p, 1) # Softmax denominator (sum of each row)
	94	+ # -- Update m_i and l_i
	95	+ alpha = tl.math.exp(m_i - m_ij) # Update factor: exp difference between old and new max
	96	+ l_i = l_i * alpha + l_ij # Update softmax denominator
	97	+ # -- Update output accumulator --
	98	+ if HEAD_DIM < 256:
	99	+ acc_ptr = acc_ptr * alpha[:, None]
	100	+ acc_ptr = tl.dot(p_cast, v, acc_ptr)
	101	+ else:
	102	+ # 1. Load current slice of accumulator
	103	+ acc = tl.load(acc_ptr + block2d_acc)
	104	+ # 2. Update in slices (split by 1/4 of BLOCK_M to avoid ub overflow)
	105	+ for i in range(4):
	106	+ # Calculate start/end rows for current slice
	107	+ offset = i * (BLOCK_M // 4)
	108	+ # Extract slice data
	109	+ acc_i = tl.extract_slice(acc, (offset, 0), (BLOCK_M // 4, HEAD_DIM), (1, 1))
	110	+ alpha_i = tl.extract_slice(alpha, [offset], [BLOCK_M // 4], [1])
	111	+ pv_i = tl.extract_slice(pv, (offset, 0), (BLOCK_M // 4, HEAD_DIM), (1, 1))
	112	+ # Incrementally update slice: acc = acc * alpha + pv
	113	+ acc_i = acc_i * alpha_i[:, None] + pv_i
	114	+ # Write updated slice back to accumulator
	115	+ acc = tl.insert_slice(acc, acc_i, (offset, 0), (BLOCK_M // 4, HEAD_DIM), (1, 1))
	116	+ # 3. updated accumulator
	117	+ tl.store(acc_ptr + block2d_acc, acc)
	118	+
	119	+ m_i = m_ij # Update current block max
	120	+ # Advance V and K block pointers to next BLOCK_N range
	121	+ V_block_ptr = tl.advance(V_block_ptr, (BLOCK_N, 0))
	122	+ K_block_ptr = tl.advance(K_block_ptr, (BLOCK_N, 0))
	123	+ # Return accumulated output acc_ptr, softmax denominator l_i, and max value m_i
	124	+ return acc_ptr, l_i, m_i
	125	+
	126	+
	127	+@triton.jit
	128	+def _attn_fwd(Q, K, V, M, Out, acc, sm_scale,
	129	+ stride_qz: tl.constexpr, stride_qh: tl.constexpr, stride_qm: tl.constexpr, stride_qk: tl.constexpr,
	130	+ stride_kz: tl.constexpr, stride_kh: tl.constexpr, stride_kn: tl.constexpr, stride_kk: tl.constexpr,
	131	+ stride_vz: tl.constexpr, stride_vh: tl.constexpr, stride_vn: tl.constexpr, stride_vk: tl.constexpr,
	132	+ stride_oz: tl.constexpr, stride_oh: tl.constexpr, stride_om: tl.constexpr, stride_on: tl.constexpr,
	133	+ Z: tl.constexpr, H: tl.constexpr,
	134	+ N_CTX: tl.constexpr,
	135	+ HEAD_DIM: tl.constexpr,
	136	+ BLOCK_M: tl.constexpr,
	137	+ BLOCK_N: tl.constexpr,
	138	+ STAGE: tl.constexpr
	139	+ ):
	140	+ # Total number of blocks in sequence dimension (M)
	141	+ NUM_BLOCKS_M = N_CTX // BLOCK_M
	142	+ # Total tasks = number of sequence blocks × batch size (Z) × number of attention heads (H)
	143	+ NUM_BLOCKS = NUM_BLOCKS_M * Z * H
	144	+
	145	+ # Current M-dimension block index
	146	+ pid = tl.program_id(0)
	147	+
	148	+ for block_idx in range(pid, NUM_BLOCKS, 20):
	149	+ task_hz_idx = block_idx // NUM_BLOCKS_M
	150	+ task_m_idx = block_idx % NUM_BLOCKS_M
	151	+ off_z = task_hz_idx // H
	152	+ off_h = task_hz_idx % H
	153	+ qvk_offset = off_z.to(tl.int64) * stride_qz + off_h.to(tl.int64) * stride_qh
	154	+ # Create block pointers for Q, K, V, Output
	155	+ Q_block_ptr = tl.make_block_ptr(
	156	+ base=Q + qvk_offset,
	157	+ shape=(N_CTX, HEAD_DIM),
	158	+ strides=(stride_qm, stride_qk),
	159	+ offsets=(task_m_idx * BLOCK_M, 0),
	160	+ block_shape=(BLOCK_M, HEAD_DIM),
	161	+ order=(1, 0),
	162	+ )
	163	+ V_block_ptr = tl.make_block_ptr(
	164	+ base=V + qvk_offset,
	165	+ shape=(N_CTX, HEAD_DIM),
	166	+ strides=(stride_vn, stride_vk),
	167	+ offsets=(0, 0),
	168	+ block_shape=(BLOCK_N, HEAD_DIM),
	169	+ order=(1, 0),
	170	+ )
	171	+ K_block_ptr = tl.make_block_ptr(
	172	+ base=K + qvk_offset,
	173	+ shape=(N_CTX, HEAD_DIM),
	174	+ strides=(stride_kn, stride_kk),
	175	+ offsets=(0, 0),
	176	+ block_shape=(BLOCK_N, HEAD_DIM),
	177	+ order=(1, 0),
	178	+ )
	179	+ O_block_ptr = tl.make_block_ptr(
	180	+ base=Out + qvk_offset,
	181	+ shape=(N_CTX, HEAD_DIM),
	182	+ strides=(stride_om, stride_on),
	183	+ offsets=(task_m_idx * BLOCK_M, 0),
	184	+ block_shape=(BLOCK_M, HEAD_DIM),
	185	+ order=(1, 0),
	186	+ )
	187	+ # Initialize offsets
	188	+ offs_m = task_m_idx * BLOCK_M + tl.arange(0, BLOCK_M)
	189	+ offs_n = tl.arange(0, BLOCK_N)
	190	+
	191	+ m_i = tl.zeros([BLOCK_M], dtype=tl.float32) - float("inf")
	192	+ l_i = tl.zeros([BLOCK_M], dtype=tl.float32) + 1.0
	193	+
	194	+ # Initialize accumulator
	195	+ if HEAD_DIM < 256:
	196	+ acc_ptr = tl.zeros([BLOCK_M, HEAD_DIM], dtype=tl.float32)
	197	+ else:
	198	+ acc_offset = (
	199	+ off_z.to(tl.int64) * stride_qz // stride_qm * HEAD_DIM +
	200	+ off_h.to(tl.int64) * stride_qh // stride_qm * HEAD_DIM +
	201	+ task_m_idx * BLOCK_M * HEAD_DIM
	202	+ )
	203	+ acc_ptr = acc + acc_offset
	204	+
	205	+ q = tl.load(Q_block_ptr)
	206	+
	207	+ # stage 1: off-band
	208	+ # For causal = True, STAGE = 3 and _attn_fwd_inner gets 1 as its STAGE
	209	+ # For causal = False, STAGE = 1, and _attn_fwd_inner gets 3 as its STAGE
	210	+ if STAGE & 1:
	211	+ acc_ptr, l_i, m_i = _attn_fwd_inner(acc_ptr, l_i, m_i, q, K_block_ptr, V_block_ptr, #
	212	+ task_m_idx, sm_scale, #
	213	+ BLOCK_M, HEAD_DIM, BLOCK_N, #
	214	+ 4 - STAGE, offs_m, offs_n, N_CTX, V.dtype.element_ty == tl.float8e5 #
	215	+ )
	216	+ # stage 2: on-band
	217	+ if STAGE & 2:
	218	+ # barrier makes it easier for compiler to schedule the
	219	+ # two loops independently
	220	+ acc_ptr, l_i, m_i = _attn_fwd_inner(acc_ptr, l_i, m_i, q, K_block_ptr, V_block_ptr, #
	221	+ task_m_idx, sm_scale, #
	222	+ BLOCK_M, HEAD_DIM, BLOCK_N, #
	223	+ 2, offs_m, offs_n, N_CTX, V.dtype.element_ty == tl.float8e5 #
	224	+ )
	225	+
	226	+ m_i += tl.math.log(l_i)
	227	+ if HEAD_DIM < 256:
	228	+ accumulator = acc_ptr / l_i[:, None]
	229	+ else:
	230	+ row = tl.arange(0, BLOCK_M)[:, None]
	231	+ col_head_dim = tl.arange(0, HEAD_DIM)[None, :]
	232	+ block2d_acc = row * HEAD_DIM + col_head_dim
	233	+ accumulator = tl.load(acc_ptr + block2d_acc)
	234	+ accumulator = accumulator / l_i[:, None]
	235	+
	236	+ m_ptrs = M + task_hz_idx * N_CTX + offs_m
	237	+
	238	+ tl.store(m_ptrs, m_i)
	239	+ tl.store(O_block_ptr, accumulator.to(Out.type.element_ty))
	240	+
	241	+
	242	+class _attention(torch.autograd.Function):
	243	+
	244	+ @staticmethod
	245	+ def forward(ctx, q, k, v, causal, sm_scale, BM, BN):
	246	+ """
	247	+ Forward computation interface:
	248	+ Args:
	249	+ ctx: Context object
	250	+ q: Query tensor (Q), shape [Z, H, N_CTX, HEAD_DIM]
	251	+ k: Key tensor (K), shape [Z, H, N_CTX, HEAD_DIM]
	252	+ v: Value tensor (V), shape [Z, H, N_CTX, HEAD_DIM]
	253	+ causal: Whether to enable causal attention
	254	+ sm_scale: Scaling factor for QK product
	255	+ BM: Q block size (BLOCK_M)
	256	+ BN: K/V block size (BLOCK_N)
	257	+ Returns:
	258	+ o: Attention output tensor, shape [Z, H, N_CTX, HEAD_DIM]
	259	+ """
	260	+ # shape constraints
	261	+ HEAD_DIM_Q, HEAD_DIM_K = q.shape[-1], k.shape[-1]
	262	+ # when v is in float8_e5m2 it is transposed.
	263	+ HEAD_DIM_V = v.shape[-1]
	264	+ assert HEAD_DIM_Q == HEAD_DIM_K and HEAD_DIM_K == HEAD_DIM_V
	265	+ assert HEAD_DIM_K in {16, 32, 64, 128, 256}
	266	+
	267	+ o = torch.empty_like(q)
	268	+ stage = 3 if causal else 1
	269	+ extra_kern_args = {}
	270	+
	271	+
	272	+ # Number of NPU cores (adjust based on hardware)
	273	+ num_cores = 20
	274	+ acc = torch.zeros((q.shape[0], q.shape[1], q.shape[2], HEAD_DIM_K), dtype=torch.float32, device=q.device)
	275	+ M = torch.empty((q.shape[0], q.shape[1], q.shape[2]), device=q.device, dtype=torch.float32)
	276	+
	277	+ _attn_fwd[(num_cores,)](
	278	+ q, k, v, M, o, acc, sm_scale,
	279	+ q.stride(0), q.stride(1), q.stride(2), q.stride(3),
	280	+ k.stride(0), k.stride(1), k.stride(2), k.stride(3),
	281	+ v.stride(0), v.stride(1), v.stride(2), v.stride(3),
	282	+ o.stride(0), o.stride(1), o.stride(2), o.stride(3),
	283	+ q.shape[0], q.shape[1], N_CTX=q.shape[2],
	284	+ HEAD_DIM=HEAD_DIM_K,
	285	+ BLOCK_M=BM,
	286	+ BLOCK_N=BN,
	287	+ STAGE=stage,
	288	+ **extra_kern_args)
	289	+
	290	+ ctx.save_for_backward(q, k, v, o, M)
	291	+ ctx.sm_scale = sm_scale
	292	+ ctx.HEAD_DIM = HEAD_DIM_K
	293	+ ctx.causal = causal
	294	+ return o
	295	+
	296	+attention = _attention.apply
	297	+
	298	+
	299	+@pytest.mark.parametrize("Z, H, N_CTX, HEAD_DIM, causal, dtype, BM, BN", [
	300	+ (1, 1, 128, 128, False, torch.float16, 32, 128),
	301	+ (1, 1, 128, 128, False, torch.bfloat16, 64, 128),
	302	+ (1, 2, 256, 256, False, torch.bfloat16, 32, 256),
	303	+ (2, 2, 128, 256, False, torch.float16, 64, 128),
	304	+ (4, 32, 64, 64, False, torch.float16, 32, 64),
	305	+ (4, 32, 1024, 64, False, torch.bfloat16, 64, 128),
	306	+ (4, 32, 4096, 64, False, torch.float16, 128, 128),
	307	+])
	308	+def test_op(Z, H, N_CTX, HEAD_DIM, causal, dtype, BM, BN):
	309	+ # Filter out non-integer cases; N_CTX must be divisible by BM and BN, and HEAD_DIM must be divisible by 16.
	310	+ if N_CTX % BM != 0 or N_CTX % BN != 0 or HEAD_DIM % 16 != 0:
	311	+ pytest.skip("Skipping non-divisible case")
	312	+
	313	+ torch.manual_seed(20)
	314	+ q = (torch.empty((Z, H, N_CTX, HEAD_DIM), dtype=dtype, device=DEVICE).normal_(mean=0.0, std=0.5).requires_grad_())
	315	+ k = (torch.empty((Z, H, N_CTX, HEAD_DIM), dtype=dtype, device=DEVICE).normal_(mean=0.0, std=0.5).requires_grad_())
	316	+ v = (torch.empty((Z, H, N_CTX, HEAD_DIM), dtype=dtype, device=DEVICE).normal_(mean=0.0, std=0.5).requires_grad_())
	317	+
	318	+ sm_scale = 0.5
	319	+
	320	+ tri_out = attention(q, k, v, causal, sm_scale, BM, BN)
	321	+ ref_out = torch_npu.npu_fusion_attention(
	322	+ q, k, v, H,
	323	+ padding_mask=None,
	324	+ atten_mask=None,
	325	+ scale=sm_scale,
	326	+ keep_prob=1.0,
	327	+ input_layout="BNSD",
	328	+ pre_tockens=65535,
	329	+ next_tockens=65535,
	330	+ sparse_mode=0,
	331	+ )[0]
	332	+
	333	+ torch.testing.assert_close(ref_out, tri_out, atol=1e-2, rtol=1e-2, equal_nan=True)
	334	+ print(f"[PASSED] Attention shape:({Z}, {H}, {N_CTX}, {HEAD_DIM}), BM: {BM}, BN: {BN}, dtype: {dtype}")
	335	+
	336	+
	337	+if __name__ == "__main__":
	338	+ test_op(1, 1, 128, 128, causal=False, dtype=torch.float16, BM=32, BN=128)
	339	+ test_op(1, 1, 128, 128, causal=False, dtype=torch.bfloat16, BM=64, BN=128)
	340	+ test_op(1, 2, 256, 256, causal=False, dtype=torch.bfloat16, BM=32, BN=256)
	341	+ test_op(2, 2, 128, 256, causal=False, dtype=torch.float16, BM=64, BN=128)
	342	+ test_op(4, 32, 64, 64, causal=False, dtype=torch.float16, BM=32, BN=64)
	343	+ test_op(4, 32, 1024, 64, causal=False, dtype=torch.bfloat16, BM=64, BN=128)
	344	+ test_op(4, 32, 4096, 64, causal=False, dtype=torch.float16, BM=128, BN=128)
	345	+```
	346	+
	347	+Output:
	348	+
	349	+```bash
	350	+[PASSED] Attention shape:(1, 1, 128, 128), BM: 32, BN: 128, dtype: torch.float16
	351	+[PASSED] Attention shape:(1, 1, 128, 128), BM: 64, BN: 128, dtype: torch.bfloat16
	352	+[PASSED] Attention shape:(1, 2, 256, 256), BM: 32, BN: 256, dtype: torch.bfloat16
	353	+[PASSED] Attention shape:(2, 2, 128, 256), BM: 64, BN: 128, dtype: torch.float16
	354	+[PASSED] Attention shape:(4, 32, 64, 64), BM: 32, BN: 64, dtype: torch.float16
	355	+[PASSED] Attention shape:(4, 32, 1024, 64), BM: 64, BN: 128, dtype: torch.bfloat16
	356	+[PASSED] Attention shape:(4, 32, 4096, 64), BM: 128, BN: 128, dtype: torch.float16
	357	+```
	358	+
	359	+The preceding logs indicate that the output on Triton is the same as that on PyTorch.

Adocs/en/examples/05_matrix_multiplication_example.md+179-0

@@ -0,0 +1,179 @@
	1	+# Matrix Multiplication
	2	+
	3	+This section describes how to implement a matrix multiplication kernel using Triton.
	4	+
	5	+## Compute Kernel
	6	+
	7	+The following Triton kernel implements batched matrix multiplication with bias:
	8	+The formula is as follows:
	9	+$$ \text{output}[b, i, j] = \sum_k \text{x}[b, i, k] \cdot \text{y}[k, j] + \text{z}[b, i, j] $$
	10	+Specifically:
	11	+- The shape of `x` is `(A, B)`.
	12	+- The shape of `y` is `(B, C)`.
	13	+- The shape of `z` (bias) is `(A, C)`.
	14	+- The shape of `output` is `(A, C)`.
	15	+
	16	+This kernel assumes that a single block is responsible for computing the entire output matrix. It is applicable to small-scale matrices (A, B, and C are small and can be fully covered by the current program block).
	17	+
	18	+```python
	19	+import pytest
	20	+import torch
	21	+import torch_npu
	22	+import triton
	23	+import triton.language as tl
	24	+
	25	+
	26	+@triton.jit
	27	+def triton_dot_2_Bias(
	28	+ output_ptr, # Pointer to the output tensor, with shape (A, C)
	29	+ x_ptr, # Pointer to the input tensor x, with shape (A, B)
	30	+ y_ptr, # Pointer to the input tensor y, with shape (B, C)
	31	+ z_ptr, # Pointer to the bias tensor z, with shape (A, C)
	32	+ A: tl.constexpr, # Size of the first dimension (batch/number of rows)
	33	+ B: tl.constexpr, # Shared dimension (number of columns in x and number of rows in y)
	34	+ C: tl.constexpr # Size of the second dimension (number of columns)
	35	+):
	36	+ # Create an index vector.
	37	+ bidx = tl.arange(0, A) # [0, 1,..., A-1], used for the row dimension.
	38	+ cidx = tl.arange(0, B) # [0, 1,..., B-1], used for the columns of x or rows of y.
	39	+ didx = tl.arange(0, C) # [0, 1,..., C-1], used for the column dimension.
	40	+
	41	+ # Construct the linear index of x: (A, B) -> flattened to A*B
	42	+ Xidx = bidx[:, None] * B + cidx[None, :] # Broadcast to form an (A, B) index grid
	43	+
	44	+ # Construct the linear index of y: (B, C) -> flattened to B*C
	45	+ Yidx = cidx[:, None] * C + didx[None, :] # (B, C) index grid
	46	+
	47	+ # Construct the linear index of z and output: (A, C).
	48	+ Zidx = bidx[:, None] * C + didx[None, :] # (A, C) index grid
	49	+
	50	+ # Load data from global memory.
	51	+ X = tl.load(x_ptr + Xidx) # Load the (A, B) sub-block.
	52	+ Y = tl.load(y_ptr + Yidx) # Load the (B, C) sub-block.
	53	+ Z = tl.load(z_ptr + Zidx) # Load the bias (A, C).
	54	+
	55	+ # Perform matrix multiplication and add the bias.
	56	+ ret = tl.dot(X, Y) + Z # tl.dot performs (A, B) × (B, C) → (A, C).
	57	+
	58	+ # Write the result back to global memory.
	59	+ oidx = bidx[:, None] * C + didx[None, :] # Same as Zidx, which can be reused.
	60	+ tl.store(output_ptr + oidx, ret)
	61	+```
	62	+## Tools and Methods
	63	+
	64	+The following helper functions are used to support the testing and verification of Triton kernels, including PyTorch reference implementation, data type mapping, random tensor generation, and result verification.
	65	+
	66	+```Python
	67	+def torch_dot_Bias(x0, x1, bias):
	68	+ """PyTorch reference implementation: Perform matrix multiplication and add the bias."""
	69	+ res = torch.matmul(x0, x1) + bias
	70	+ return res
	71	+
	72	+def get_torch_typename(dtype):
	73	+ """Map the data type in string format to the corresponding torch.dtype."""
	74	+ if dtype == 'float32':
	75	+ tyname = torch.float32
	76	+ elif dtype == 'int32':
	77	+ tyname = torch.int32
	78	+ elif dtype == 'int64':
	79	+ tyname = torch.int64
	80	+ elif dtype == 'float16':
	81	+ tyname = torch.float16
	82	+ elif dtype == 'int16':
	83	+ tyname = torch.int16
	84	+ elif dtype == 'int8':
	85	+ tyname = torch.int8
	86	+ elif dtype == 'bool':
	87	+ tyname = torch.bool
	88	+ elif dtype == 'bfloat16':
	89	+ tyname = torch.bfloat16
	90	+ else:
	91	+ raise ValueError('Invalid parameter \"dtype\" is found : {}'.format(dtype))
	92	+ return tyname
	93	+
	94	+def generate_tensor(shape, dtype):
	95	+ """Generates a random tensor based on the specified shape and data type, and adapts to the value ranges of different data types."""
	96	+ if dtype == 'float32' or dtype == 'float16' or dtype == 'bfloat16':
	97	+ return torch.randn(size=shape, dtype=eval('torch.' + dtype))
	98	+ elif dtype == 'int32' or dtype == 'int64' or dtype == 'int16':
	99	+ return torch.randint(low=0, high=2000, size=shape, dtype=eval('torch.' + dtype))
	100	+ elif dtype == 'int8':
	101	+ return torch.randint(low=0, high=127, size=shape, dtype=eval('torch.' + dtype))
	102	+ elif dtype == 'bool':
	103	+ return torch.randint(low=0, high=2, size=shape).bool()
	104	+ else:
	105	+ raise ValueError('Invalid parameter \"dtype\" is found : {}'.format(dtype))
	106	+
	107	+def validate_cmp(dtype, y_cal, y_ref):
	108	+ """Compare the Triton compute result with the PyTorch reference result on the NPU, and set the tolerance or strict equality based on the data type."""
	109	+ y_cal=y_cal.npu()
	110	+ y_ref=y_ref.npu()
	111	+ if dtype == 'float16':
	112	+ torch.testing.assert_close(y_ref, y_cal, rtol=1e-03, atol=1e-03, equal_nan=True)
	113	+ elif dtype == 'bfloat16':
	114	+ torch.testing.assert_close(y_ref.to(torch.float32), y_cal.to(torch.float32), rtol=1e-03, atol=1e-03, equal_nan=True)
	115	+ elif dtype == 'float32':
	116	+ torch.testing.assert_close(y_ref, y_cal, rtol=1e-04, atol=1e-04, equal_nan=True)
	117	+ elif dtype == 'int32' or dtype == 'int64' or dtype == 'int16' or dtype == 'int8':
	118	+ assert torch.equal(y_cal, y_ref)
	119	+ elif dtype == 'bool':
	120	+ assert torch.equal(y_cal, y_ref)
	121	+ else:
	122	+ raise ValueError('Invalid parameter \"dtype\" is found : {}'.format(dtype))
	123	+```
	124	+
	125	+## Parameterized Test
	126	+
	127	+Use `pytest` to verify the parameterization function of the `triton_dot_2_Bias` kernel, covering different combinations of matrix dimensions and data types.
	128	+
	129	+```python
	130	+# Test case configuration: (A, B, C) indicates that matrix x is (A, B), y is (B, C), and bias/output is (A, C).
	131	+testlist = [
	132	+ (16, 16, 16),
	133	+]
	134	+
	135	+# Supported data types (only float16 is supported currently)
	136	+typelist = ['float16',]
	137	+
	138	+@pytest.mark.parametrize('A, B, C', testlist)
	139	+@pytest.mark.parametrize('sigtype', typelist)
	140	+def test_dot_2_Bias(sigtype, A, B, C):
	141	+ """Perform an end-to-end function test on the triton_dot_2_Bias kernel."""
	142	+ dtype = get_torch_typename(sigtype)
	143	+
	144	+ # Generate the input tensor and move it to the NPU.
	145	+ x0 = generate_tensor(shape=(A, B), dtype=sigtype).npu()
	146	+ x1 = generate_tensor(shape=(B, C), dtype=sigtype).npu()
	147	+
	148	+ # The bias items are generated using float32 (to avoid accuracy issues caused by integer bias).
	149	+ if 'int' in sigtype:
	150	+ bias = generate_tensor(shape=(A, C), dtype='int32').npu()
	151	+ # The integer input needs to be converted to float32 for computation and then converted back to the target type.
	152	+ ans = torch_dot_Bias(x0.to(torch.float32), x1.to(torch.float32), bias.to(torch.float32)).to(dtype)
	153	+ else:
	154	+ bias = generate_tensor(shape=(A, C), dtype='float32').npu()
	155	+ ans = torch_dot_Bias(x0, x1, bias).to(eval(f"torch.{dtype}"))
	156	+
	157	+ # Initialize the output tensor.
	158	+ output = torch.zeros((A, C), dtype=dtype).npu()
	159	+
	160	+ # Start the Triton kernel (grid=(1,1,1), single-block execution).
	161	+ triton_dot_2_Bias[1, 1, 1](output, x0, x1, bias, A, B, C, debug=True)
	162	+
	163	+ # Verify the result correctness.
	164	+ validate_cmp(sigtype, output, ans)
	165	+ print(f"Test matmul with dtype={sigtype}, shape=({A},{B},{C}) PASSED!")
	166	+
	167	+
	168	+if __name__ == "__main__":
	169	+ # Running a single test case directly (for debugging) is supported.
	170	+ test_dot_2_Bias("float16", 16, 16, 16)
	171	+```
	172	+
	173	+Output example:
	174	+
	175	+```python
	176	+Test matmul with dtype=float16, shape=(16,16,16) PASSED!
	177	+```
	178	+
	179	+The preceding logs indicate that the output on Triton is the same as that on PyTorch.

Adocs/en/examples/06_autotune_example.md+287-0

@@ -0,0 +1,287 @@
	1	+# Autotune
	2	+
	3	+In this section, we will demonstrate how to use the autotune method of Triton to automatically select the optimal kernel configuration parameters. Currently, Triton-Ascend autotune is fully compatible with the usage of the autotune in the community (visit https://triton-lang.org/main/python-api/generated/triton.autotune.html). That is, users need to manually pass some defined triton.Config to autotune, and then autotune selects the optimal kernel configuration through benchmarking. In addition, Triton-Ascend provides the advanced autotune usage. Users need to provide information such as the split and tiling axes of the current Triton kernel. In this case, autotune automatically generates some possible optimal kernel configurations based on the actual input size, and then selects the optimal configuration through benchmarking or profiling.
	4	+
	5	+Note:
	6	+Currently, Triton-Ascend autotune supports block size and multibuffer (compiler optimization). However, the num_warps and num_stages parameters are not supported due to hardware architecture differences. In the future, more adjustable autotune options will be added.
	7	+
	8	+## Community Autotune Usage Example
	9	+```Python
	10	+import torch, torch_npu
	11	+import triton
	12	+import triton.language as tl
	13	+
	14	+def test_triton_autotune():
	15	+
	16	+ # Return a group of different kernel configurations for autotune testing.
	17	+ def get_autotune_config():
	18	+ return [
	19	+ triton.Config({'XS': 1 * 128, 'multibuffer': True}),
	20	+ triton.Config({'XS': 12 * 1024, 'multibuffer': True}),
	21	+ triton.Config({'XS': 12 * 1024, 'multibuffer': False}),
	22	+ triton.Config({'XS': 8 * 1024, 'multibuffer': True}),
	23	+ ]
	24	+
	25	+ @triton.autotune(
	26	+ configs=get_autotune_config(), # Configuration list
	27	+ key=["numel"], # Autotune is triggered when the numel size changes.
	28	+ )
	29	+ @triton.jit
	30	+ def triton_calc_kernel(
	31	+ out_ptr0, in_ptr0, in_ptr1, numel,
	32	+ XS: tl.constexpr # Block size, which is used to control the amount of data processed by each thread block.
	33	+ ):
	34	+ pid = tl.program_id(0) # Obtain the ID of the current program.
	35	+ idx = pid * XS + tl.arange(0, XS) # Index range processed by the current thread block.
	36	+ msk = idx < numel # Mask to avoid out-of-bounds access.
	37	+
	38	+ # Repeat computation to simulate load (for perf test).
	39	+ for i in range(10000):
	40	+ tmp0 = tl.load(in_ptr0 + idx, mask=msk, other=0.0) # Load x0.
	41	+ tmp1 = tl.load(in_ptr1 + idx, mask=msk, other=0.0) # Load x1.
	42	+ tmp2 = tl.math.exp(tmp0) + tmp1 + i # Compute.
	43	+ tl.store(out_ptr0 + idx, tmp2, mask=msk) # Store and output the result.
	44	+
	45	+ # Triton calls a function and automatically uses the autotuned kernel.
	46	+ def triton_calc_func(x0, x1):
	47	+ n = x0.numel()
	48	+ y0 = torch.empty_like(x0)
	49	+ grid = lambda meta: (triton.cdiv(n, meta["XS"]), 1, 1) # Compute the grid size.
	50	+ triton_calc_kernel[grid](y0, x0, x1, n)
	51	+ return y0
	52	+
	53	+ # Use PyTorch as the reference implementation for comparison.
	54	+ def torch_calc_func(x0, x1):
	55	+ return torch.exp(x0) + x1 + 10000 - 1
	56	+
	57	+ DEV = "npu" # Use the NPU as the device.
	58	+ DTYPE = torch.float32
	59	+ N = 192 * 1024 # Input length.
	60	+ x0 = torch.randn((N,), dtype=DTYPE, device=DEV) # Randomly input x0.
	61	+ x1 = torch.randn((N,), dtype=DTYPE, device=DEV) # Randomly input x1.
	62	+ torch_ref = torch_calc_func(x0, x1) # Obtain the reference result.
	63	+ triton_cal = triton_calc_func(x0, x1) # Run the Triton kernel.
	64	+ torch.testing.assert_close(triton_cal, torch_ref) # Verify whether the outputs are consistent.
	65	+
	66	+if __name__ == "__main__":
	67	+ test_triton_autotune()
	68	+ print("success: test_triton_autotune") # Print success message.
	69	+```
	70	+
	71	+## Advanced Autotune Usage Example
	72	+```Python
	73	+# The following are parameters added or modified compared with the community autotune.
	74	+# Note: When either split_params or tiling_params is not empty, the advanced autotune method is automatically triggered.
	75	+
	76	+# In the dictionary consisting of "key (Dict[str, str]): axis name: argument name", the change of the argument triggers the regeneration and evaluation of candidate configurations.
	77	+# The axis name belongs to the set {'x', 'y', 'z', 'w', 'v', 't', 'rx', 'ry', 'rz', 'rw', 'rv', 'rt'}. The prefix 'r' indicates the reduction axis.
	78	+# The prefix 'r' should be added only when the axis name in this parameter is used as the reduction axis.
	79	+# In the dictionary consisting of "split_params (Dict[str, str]): axis name: argument name", the argument is the tunable parameter of the split axis, for example, 'XBLOCK'.
	80	+# The axis name must be in the axis name set of the parameter `key`. Do not prefix the axis name with 'r'.
	81	+# This parameter can be left empty. If both split_params and tiling_params are empty, autotune is not performed.
	82	+# The split axis can be determined based on the kernel splitting statement `tl.program_id()`.
	83	+# In the dictionary consisting of "tiling_params (Dict[str, str]): axis name: argument name", the argument is an tunable parameter of the tiling axis, for example, 'XBLOCK_SUB'.
	84	+# The axis name must be in the axis name set of the parameter `key`. Do not prefix the axis name with 'r'.
	85	+# This parameter can be left empty. If both split_params and tiling_params are empty, autotune is not performed.
	86	+# The tiling axis can be determined based on the `tl.arange()` expression.
	87	+# low_dims (List[str]): list of axis names of all low-dimensional axes. The axis name must be in the axis name set of the parameter `key`. Do not prefix the axis name with 'r'.
	88	+# dual_reduction (bool): specifies whether to perform reduction on multiple axes, which affects the tiling generation policy.
	89	+# persistent_reduction (bool): specifies whether to perform tiling on the reduction axis, which affects the tiling generation policy.
	90	+# For details, see the cases in ascend\examples\autotune_cases.
	91	+@triton.autotune(
	92	+ configs=[],
	93	+ key={"x": "n_elements"}, # Size of the split axis x.
	94	+ split_params={"x": "BLOCK_SIZE"}, # Size of BLOCK_SIZE to be adjusted for the split axis x.
	95	+ tiling_params={}, # The tiling axis is the split axis.
	96	+ low_dims=["x"], # Low-dimensional axis.
	97	+ persistent_reduction=False,
	98	+ dual_reduction=False,
	99	+)
	100	+@triton.jit
	101	+def add_kernel(
	102	+ x_ptr, # Pointer to the first input vector.
	103	+ y_ptr, # Pointer to the second input vector.
	104	+ output_ptr, # Pointer to the output vector.
	105	+ n_elements, # Size of the vector.
	106	+ BLOCK_SIZE: tl.constexpr, # Number of elements that should be processed by each kernel.
	107	+ # Note: `constexpr` indicates that it can be determined at compile time and therefore can be used as a shape value.
	108	+):
	109	+ pid = tl.program_id(axis=0) # A one-dimensional grid is used, so the axis is 0.
	110	+ # Offset of the data to be processed by the current kernel in the memory relative to the start address.
	111	+ # For example, if there is a vector of length 256 and block sizes 64, each program
	112	+ # will access the elements [0:64, 64:128, 128:192, 192:256] respectively.
	113	+ # Note that offsets is a list of pointers:
	114	+ block_start = pid * BLOCK_SIZE
	115	+ offsets = block_start + tl.arange(0, BLOCK_SIZE)
	116	+ # Create a mask to prevent out-of-bounds memory access.
	117	+ mask = offsets < n_elements
	118	+ # Load x and y, and use the mask to mask out the redundant elements to prevent the length of the input vector from not being an integer multiple of the block size.
	119	+ x = tl.load(x_ptr + offsets, mask=mask)
	120	+ y = tl.load(y_ptr + offsets, mask=mask)
	121	+ output = x + y
	122	+ # Write x + y back.
	123	+ tl.store(output_ptr + offsets, output, mask=mask)
	124	+```
	125	+
	126	+Note:
	127	+1. By default, Triton-Ascend uses the benchmark mode to obtain the on-chip computation time. After the environment variable is set by running `export TRITON_BENCH_METHOD="npu"`, the on-chip computation time of each kernel is obtained by using `torch_npu.profiler.profile`. For some Triton kernels that compute fast, such as small-shape operators, this method can obtain more accurate computation time than the default method. However, this will significantly increase the overall autotune time. Therefore, exercise caution when enabling this method.
	128	+2. Currently, this advanced usage is mainly used for vector operators and is not supported by cube operators. For more advanced usage examples, see [Advanced Autotune Cases](https://gitcode.com/Ascend/triton-ascend/tree/master/ascend/examples/autotune_cases).
	129	+
	130	+### Automatic Parameter Parsing
	131	+
	132	+Before automatically parsing parameters, the system obtains the parameters that are not passed during the `kernel` function call. The parameters that are not passed are used as the candidate parameters for the split axis and tiling axis.
	133	+
	134	+```Python
	135	+@triton.jit
	136	+def kernel_func(
	137	+ outputptr,
	138	+ input_ptr,
	139	+ n_rows,
	140	+ n_cols,
	141	+ BLOCK_SIZE: tl.constexpr,
	142	+ XBLOCK: tl.constexpr,
	143	+ XBLOCK_SUB: tl.constexpr,
	144	+):
	145	+ # kernel implementation
	146	+ ...
	147	+
	148	+# If XBLOCK and XBLOCK_SUB are not passed, they are used as candidate parameters for the split axis and tiling axis.
	149	+# BLOCK_SIZE is passed as a keyword argument and is not used as a candidate parameter. Therefore, it will not be identified.
	150	+kernel_func[grid](y, x, n_rows, n_cols, BLOCK_SIZE=block_size)
	151	+```
	152	+
	153	+#### Split Axis Parameter Parsing
	154	+
	155	+The split axis parameters are parsed based on the kernel splitting statement `tl.program_id()`. The system analyzes the usage of the `tl.program_id()` variable in the program and the multiplication operation between the variable and other variables to identify potential split axis parameters (currently, direct or indirect multiplication through intermediate variables is supported) and filters the parameters based on the candidate parameter list (parameters not provided by users).
	156	+
	157	+Finally, the split axis corresponding to the current parameters is identified through mask comparison and the `key` passed in `autotune`.
	158	+
	159	+Notes: 1. The split axis parameter must be multiplied by `tl.program_id()`. 2. The mask comparison must be performed, and the `key` corresponding to the split axis or the min function with the `key` as the parameter must be used as the right value. Otherwise, the axis cannot be identified and the parameter parsing will fail.3. The identified axis parameters are limited to the candidate parameter list. This ensures that only the parameters that can be dynamically tuned by autotune are considered.
	160	+
	161	+```Python
	162	+@triton.autotune(
	163	+ key={"n_elements"} # It needs to be specified.
	164	+ ...
	165	+)
	166	+@triton.jit
	167	+def triton_func(...):
	168	+ # case1:
	169	+ pid = tl.program_id(0)
	170	+ block_start = pid * XBLOCK
	171	+ offsets = block_start + tl.arange(0, XBLOCK)
	172	+
	173	+ # case2:
	174	+ block_start = tl.program_id(0) * XBLOCK
	175	+ offsets = block_start + tl.arange(0, XBLOCK)
	176	+
	177	+ # case3:
	178	+ offsets = tl.program_id(0) * XBLOCK + tl.arange(0, XBLOCK)
	179	+
	180	+ # mask compare
	181	+ mask = offsets < n_elements # 1
	182	+ mask = offsets < min(..., n_elements) # 2
	183	+
	184	+# The split axis parameter split_params is parsed as {"x": "XBLOCK"}.
	185	+```
	186	+
	187	+#### Tiling Axis Parameter Parsing
	188	+
	189	+The tiling axis parameter is determined based on the `tl.arange()`, `tl.range()`, and `range()` tiling statements. The potential tiling axis parameters are identified by analyzing the usage of `tl.range()`, `tl.arange()`, and `range()` in the `for` loop in the program, and the variables computed based on the usage. The common parameters of `tl.range()` or `range()` and `tl.arange()` are extracted and filtered based on the candidate parameter list (parameters not provided by users).
	190	+
	191	+Finally, the split axis corresponding to the current parameter is identified through mask comparison with the `key` passed in `autotune`.
	192	+
	193	+Notes: 1. The tiling axis parameters must be used in the call of `tl.arange()` and be involved in the computation of the loop range in the `for` loop through `tl.range()`, `range()`, or integer division (`//`). 2. The mask comparison must be performed, and the key corresponding to the tiling axis or the min function with the key as the parameter must be used as the right value. Otherwise, the axis cannot be identified and the parameter parsing will fail.3. The identified tiling parameters are limited to the candidate parameter list. This ensures that only the parameters that can be dynamically tuned by autotune are considered.
	194	+
	195	+```Python
	196	+@triton.autotune(
	197	+ key={"n_rows", "n_cols"} # It needs to be specified.
	198	+ ...
	199	+)
	200	+@triton.jit
	201	+def triton_func(...):
	202	+ ...
	203	+ # case 1
	204	+ for row_idx in tl.range(0, XBLOCK, XBLOCK_SUB):
	205	+ row_offsets = row_idx + tl.arange(0, XBLOCK_SUB)[:, None]
	206	+ col_offsets = tl.arange(0, BLOCK_SIZE)[None, :]
	207	+
	208	+ # case 2
	209	+ loops = (XBLOCK + XBLOCK_SUB - 1) // XBLOCK_SUB
	210	+ for loop in range(loops):
	211	+ row_offsets = loop * XBLOCK_SUB + tl.arange(0, XBLOCK_SUB)[:, None]
	212	+ col_offsets = tl.arange(0, BLOCK_SIZE)[None, :]
	213	+
	214	+ ...
	215	+ xmask = row_offsets < n_rows # 1
	216	+ xmask = row_offsets < min(..., n_rows) # 2
	217	+ ymask = col_offsets < n_cols
	218	+
	219	+# The tiling axis parameter tiling_params is parsed as {"x": "XBLOCK_SUB"}.
	220	+# Although the BLOCK_SIZE parameter is also in tl.arange and is compared with n_cols to compute the mask, it is not a tiling axis parameter.
	221	+```
	222	+
	223	+#### Low-Dimensional Axis Parameter Parsing
	224	+
	225	+The low-dimensional axis parameters are parsed based on the tiling statement `tl.arange()`. The potential low-dimensional axis parameters are identified by analyzing the usage of `tl.arange()` in the program and the variables computed by it. `tl.arange()` and the variables involved in the computation are extracted. The dimension is expanded based on whether slicing is performed, and the filtering is performed based on the expansion of the dimension.
	226	+
	227	+Finally, the low-dimensional axis of the current kernel is determined by comparing the mask with the `key` passed in `autotune`.
	228	+
	229	+Notes: 1. The low-dimensional axis must be computed using `tl.arange()` and sliced. It will be identified only when expansion is perform on or slicing is not involved in the non-lowest dimension. 2. If mask comparison is not performed, the specific low-dimensional axis cannot be identified, resulting in parameter parsing failure.
	230	+
	231	+```Python
	232	+@triton.autotune(
	233	+ key={"n_rows", "n_cols"} # Automatically allocated in the order of {"x": "n_rows", "y": "n_cols"}
	234	+ ...
	235	+)
	236	+@triton.jit
	237	+def triton_func(...):
	238	+ ...
	239	+ for row_idx in tl.range(0, XBLOCK, XBLOCK_SUB):
	240	+ row_offsets = row_idx + tl.arange(0, XBLOCK_SUB)[:, None]
	241	+ col_offsets = tl.arange(0, BLOCK_SIZE)[None, :]
	242	+
	243	+ xmask = row_offsets < n_rows
	244	+ ymask = col_offsets < n_cols
	245	+
	246	+# The low-dimensional axis low_dims is parsed as {"y"}.
	247	+# Although row_offsets is also computed using tl.arange and compared with n_rows to compute the mask, slices are expanded in a low dimension. Therefore, x is not a low-dimensional axis.
	248	+```
	249	+
	250	+#### Parameter Pointer Parsing
	251	+
	252	+The pointer-type parameters are parsed based on whether the parameters are involved in the memory access statements of `tl.load()` and `tl.store()`.
	253	+
	254	+First, all parameters in the kernel function are parsed, and then all variables involved in the computation of each parameter are recursively searched.
	255	+
	256	+If a parameter is directly or indirectly (via the intermediate variable obtained by the parameter through computation) involved in the computation of the first parameter of `tl.load()` and `tl.store()`, the parameter is considered as a pointer-type parameter.
	257	+
	258	+Notes: 1. Variables modified by `tl.constexpr` are not pointer-type variables and will not be parsed subsequently. 2. Only memory access statements with directly or indirectly (via the intermediate variable obtained by parameters through computation) involved parameters are counted. If the intermediate variables obtained by these parameters are involved in the computation for more than two times, the intermediate variables are not counted.
	259	+
	260	+```Python
	261	+@triton.autotune(...)
	262	+@triton.jit
	263	+def triton_func(input_ptr, output_ptr, ...):
	264	+ ...
	265	+ # case1
	266	+ input = tl.load(input_ptr + offsets, mask=mask)
	267	+ tl.store(output_ptr + offsets, input, mask=mask)
	268	+
	269	+ # case2
	270	+ inputs_ptr = input_ptr + offsets
	271	+ input = tl.load(inputs_ptr, mask=mask)
	272	+ outputs_ptr = output_ptr + offsets
	273	+ tl.store(outputs_ptr, input, mask=mask)
	274	+
	275	+# The parsed pointer parameters are input_ptr and output_ptr.
	276	+```
	277	+
	278	+## More Functions
	279	+### Automatically Generating the Profiling Result of the Optimal Configuration
	280	+```Python
	281	+# Automatically generate the profiling result of the optimal kernel configuration of the current autotune in the `auto_profile_dir` directory, that is, the performance data collected by `torch_npu.profiler.profile`.
	282	+# This takes effect in both the community autotune usage and advanced autotune usage.
	283	+@triton.autotune(
	284	+ auto_profile_dir="./profile_result",
	285	+ ...
	286	+)
	287	+```

Adocs/en/examples/07_accuracy_comparison_example.md+119-0

@@ -0,0 +1,119 @@
	1	+# Accuracy Comparison
	2	+
	3	+In this section, you will use Triton to write a simple accuracy comparison program.
	4	+During this process, you will learn:
	5	+
	6	+- The method of comparing the accuracy of each data type in Triton.
	7	+- Reference code: triton-ascend/ascend/examples/tutorials/14-accuracy-comparison.py
	8	+
	9	+Compute kernel:
	10	+
	11	+```Python
	12	+def test_add(x0, x1):
	13	+ """
	14	+ Test the vector addition implemented by Triton and compare its accuracy with that of PyTorch.
	15	+
	16	+ Procedure:
	17	+ 1. Use PyTorch to compute the reference result (torch_ref).
	18	+ 2. Use Triton to compile the kernel and compute the result (triton_cal).
	19	+ 3. Call accuracy_comparison to compare the accuracy.
	20	+ """
	21	+
	22	+ # 1. Use PyTorch as the reference implementation (golden truth).
	23	+ def torch_func(x0, x1):
	24	+ res = x0 + x1
	25	+ return res
	26	+
	27	+ # 2. Define the Triton kernel (executed on the NPU or GPU).
	28	+ @triton.jit
	29	+ def triton_kernel_add(
	30	+ out_ptr0, # Pointer to the output: location where the result is stored
	31	+ in_ptr0, # Pointer 0 to the input: start address of x0
	32	+ in_ptr1, # Pointer 1 to the input: start address of x1
	33	+ XS: tl.constexpr # constexpr parameter: vector length, which is determined at compile time
	34	+ ):
	35	+ # Generate an index array of [0, 1, 2,..., XS-1].
	36	+ idx = tl.arange(0, XS)
	37	+ # Load the value of x0 from in_ptr0 + idx.
	38	+ tmp0 = tl.load(in_ptr0 + idx)
	39	+ # Load the value of x1 from in_ptr1 + idx.
	40	+ tmp1 = tl.load(in_ptr1 + idx)
	41	+ # Perform addition.
	42	+ tmp2 = tmp0 + tmp1
	43	+ # Write the result to out_ptr0 + idx.
	44	+ tl.store(out_ptr0 + idx, tmp2)
	45	+
	46	+ # 3. Triton encapsulation function: Call the kernel and return the result.
	47	+ def triton_func(x0, x1):
	48	+ y0 = torch.empty_like(x0) # Create an output tensor with the same shape and dtype as the input.
	49	+ # Start the kernel. grid = [1, 1, 1] indicates that only one block is used.
	50	+ # Note: XS must be passed as a parameter because it is of the tl.constexpr type.
	51	+ triton_kernel_add[1, 1, 1](y0, x0, x1, XS=x0.numel())
	52	+ return y0
	53	+
	54	+ # 4. Obtain the reference result and Triton computation result.
	55	+ torch_ref = torch_func(x0, x1)
	56	+ triton_cal = triton_func(x0, x1)
	57	+
	58	+ # 5. Compare the accuracy.
	59	+ accuracy_comparison(triton_cal, torch_ref)
	60	+
	61	+ # 6. Print success information.
	62	+ print(f"== dtype:{triton_cal.dtype} == The accuracy comparison between triton_cal and torch_ref was successful.")
	63	+
	64	+
	65	+```
	66	+
	67	+Create an accuracy comparison function that adapts to each dtype and uses the corresponding accuracy comparison method.
	68	+
	69	+```Python
	70	+
	71	+def accuracy_comparison(y_cal, y_ref):
	72	+ """
	73	+ Accuracy comparison function: Select a proper comparison policy based on the data type.
	74	+
	75	+ Processing policies for different data types:
	76	+ - Floating-point types (float16/32, bfloat16): Use torch.testing.assert_close and set the relative/absolute error tolerance.
	77	+ - Integer types (int8/16/32/64): The results must be equal (torch.equal).
	78	+ - Boolean type (bool): Strict comparison is performed on the CPU (to avoid device differences).
	79	+ """
	80	+ # Check whether the output data types are consistent.
	81	+ assert y_cal.dtype == y_ref.dtype, f"dtype mismatch: {y_cal.dtype} vs {y_ref.dtype}"
	82	+ tensor_dtype = y_cal.dtype
	83	+
	84	+ # Move the tensor to the NPU (assuming that the test is performed on the NPU).
	85	+ y_cal = y_cal.npu()
	86	+ y_ref = y_ref.npu()
	87	+
	88	+ # Select different comparison methods based on the data types.
	89	+ if tensor_dtype == torch.float16:
	90	+ # For the float16 type, the accuracy is low, and a slightly larger error is allowed.
	91	+ torch.testing.assert_close(y_ref, y_cal, rtol=1e-3, atol=1e-3, equal_nan=True)
	92	+ elif tensor_dtype == torch.bfloat16:
	93	+ # For bfloat16, the accuracy is lower. You are advised to convert it to float32 before comparison.
	94	+ torch.testing.assert_close(
	95	+ y_ref.to(torch.float32),
	96	+ y_cal.to(torch.float32),
	97	+ rtol=1e-3,
	98	+ atol=1e-3,
	99	+ equal_nan=True
	100	+ )
	101	+ elif tensor_dtype == torch.float32:
	102	+ # For the float32 type, the accuracy is high. A stricter tolerance is recommended.
	103	+ torch.testing.assert_close(y_ref, y_cal, rtol=1e-4, atol=1e-4, equal_nan=True)
	104	+ elif tensor_dtype in [torch.int64, torch.int32, torch.int16, torch.int8]:
	105	+ # For the integer type, the results must be equal.
	106	+ assert torch.equal(y_cal, y_ref), f"Integer tensors are not equal for dtype {tensor_dtype}"
	107	+ elif tensor_dtype == torch.bool:
	108	+ # For the Boolean type, comparison on the CPU is recommended to avoid differences in Boolean representation between devices.
	109	+ assert torch.equal(y_cal.cpu(), y_ref.cpu()), "Boolean tensors are not equal"
	110	+ else:
	111	+ raise ValueError(f'Invalid or unsupported tensor dtype: {tensor_dtype}')
	112	+
	113	+
	114	+```
	115	+
	116	+You can run the following command to execute the sample code: tutorials/14-accuracy-comparison.py.
	117	+```Python
	118	+python triton-ascend/ascend/examples/tutorials/14-accuracy-comparison.py
	119	+```

Adocs/en/examples/index.md+12-0

@@ -0,0 +1,12 @@
	1	+# Examples
	2	+
	3	+## Triton Examples
	4	+\|Example\|Description\|
	5	+\|--\|--\|
	6	+\| [01_vector_add_example](./01_vector_add_example.md)\| Example of simple addition\|
	7	+\| [02_fused_softmax_example](./02_fused_softmax_example.md)\| Example of Softmax fused operator\|
	8	+\| [03_layer_norm_example](./03_layer_norm_example.md)\| Example of Layer Normalization\|
	9	+\| [04_fused_attention_example](./04_fused_attention_example.md)\| Example of Flash Attention v2 fused attention algorithm\|
	10	+\| [05_matrix_multiplication_example](./05_matrix_multiplication_example.md)\| Example of efficient matrix multiplication\|
	11	+\| [06_autotune_example](./06_autotune_example.md)\| Example of using Autotune on kernels\|
	12	+\| [07_accuracy_comparison_example](./07_accuracy_comparison_example.md)\| Accuracy comparison example\|

Adocs/en/figures/FA_Performance.png+0-0

Adocs/en/figures/MM_Performance.png+0-0

Adocs/en/figures/Softmax_Performance.png+0-0

Adocs/en/figures/architectural_diagram.png+0-0

Adocs/en/figures/optimization.png+0-0

Adocs/en/figures/performance_analysis_analyse_data_code_mapping.png+0-0

Adocs/en/figures/performance_analysis_analyse_data_op_summary.png+0-0

Adocs/en/figures/performance_analysis_analyse_data_waveform.png+0-0

Adocs/en/figures/time_consumed.png+0-0

Adocs/en/figures/trace_json_with_chrome.png+0-0

Adocs/en/figures/trace_json_with_insight.png+0-0

Adocs/en/figures/visualize_data_with_insight.png+0-0

Adocs/en/installation_guide.md+282-0

@@ -0,0 +1,282 @@
	1	+# Installation Guide
	2	+
	3	+## Preparing the Environment
	4	+
	5	+### Python Version Requirements
	6	+
	7	+Triton-Ascend requires Python 3.9 to 3.11.
	8	+
	9	+### Installing Ascend CANN
	10	+
	11	+Compute Architecture for Neural Networks (CANN) is a heterogeneous compute architecture developed by Ascend for AI scenarios.
	12	+It plays a pivotal bridging role: providing upward integration with multiple AI frameworks (including MindSpore, PyTorch, and TensorFlow), while offering downward support for AI processors and programming. This establishes it as a key platform for improving the computing efficiency of Ascend AI processors.
	13	+
	14	+You can visit the Ascend community website, and install and configure CANN according to the provided software installation guide.
	15	+
	16	+During the installation, select one of the following CANN versions in {version}:
	17	+
	18	+CANN version:
	19	+
	20	+- Commercial edition
	21	+
	22	+\| Triton-Ascend Version\| CANN Commercial Version\| CANN Release Date\|
	23	+\|-------------------\|----------------------\|--------------------\|
	24	+\| 3.2.0 \| CANN 8.5.0 \| 2026-01-16 \|
	25	+\| 3.2.0rc4 \| CANN 8.3.RC2 \| 2025-11-20 \|
	26	+\| \| CANN 8.3.RC1 \| 2025-10-30 \|
	27	+
	28	+- Community edition
	29	+
	30	+\| Triton-Ascend Version\| CANN Community Version\| CANN Release Date\|
	31	+\|-------------------\|----------------------\|--------------------\|
	32	+\| 3.2.0 \| CANN 8.5.0 \| 2026-01-16 \|
	33	+\| 3.2.0rc4 \| CANN 8.3.RC2 \| 2025-11-20 \|
	34	+\| \| CANN 8.5.0.alpha001 \| 2025-11-12 \|
	35	+\| \| CANN 8.3.RC1 \| 2025-10-30 \|
	36	+
	37	+Specify the actual CPU architecture in {arch} (aarch64 or x86_64) and the software package corresponding to the software version ({version}).
	38	+
	39	+It is advisable to download and install version 8.5.0.
	40	+
	41	+\| Software Type \| Software Package Description \| Software Package Name \|
	42	+\|---------\|------------------\|----------------------------------\|
	43	+\| Toolkit \| CANN development kit \| Ascend-cann-toolkit_{version}_linux-{arch}.run \|
	44	+\| Ops \| CANN binary operator package\| Ascend-cann-A3-ops_{version}_linux-{arch}.run\|
	45	+
	46	+Note 1: The naming of A2 series Ops packages is slightly different from that of the A3 series. Here is a reference naming format: Ascend-cann-910b-ops_{version}_linux-{arch}.run.
	47	+
	48	+Note 2: The naming of Ops packages corresponding to versions earlier than 8.5.0 is slightly different. Here is a reference naming format: Atlas-A3-cann-kernels_{version}_linux-{arch}.run.
	49	+
	50	+You can find the relevant software packages at the [community](https://www.hiascend.com/developer/download/community/result?module=cann).
	51	+
	52	+The [community installation guide](https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/850/softwareinst/instg/instg_quick.html?Mode=PmIns&InstallType=local&OS=Ubuntu&Software=cannToolKit) provides the complete installation process and dependency configuration suggestions, and is therefore applicable to users opting for a comprehensive deployment of the CANN environment.
	53	+
	54	+#### CANN Installation Script
	55	+
	56	+The following uses A3 CANN 8.5.0 as an example and provides a script-based installation process for your reference.
	57	+```bash
	58	+
	59	+# Modify the execute permission on the .run packages.
	60	+chmod +x Ascend-cann-toolkit_8.5.0_linux-aarch64.run
	61	+chmod +x Ascend-cann-A3-ops_8.5.0_linux-aarch64.run
	62	+
	63	+# Common installation (default installation path: /usr/local/Ascend)
	64	+sudo ./Ascend-cann-toolkit_8.5.0_linux-aarch64.run --install
	65	+# Default installation path (same as the toolkit package: /usr/local/Ascend)
	66	+sudo ./Ascend-cann-A3-ops_8.5.0_linux-aarch64.run --install
	67	+# Put the environment variable defining the default path into effect.
	68	+source /usr/local/Ascend/ascend-toolkit/set_env.sh
	69	+
	70	+# Install the Python dependencies of CANN.
	71	+pip install attrs==24.2.0 numpy==1.26.4 scipy==1.13.1 decorator==5.1.1 psutil==6.0.0 pyyaml
	72	+```
	73	+
	74	+- Note: If the installation path is not specified, software will be installed in the default path. The default installation paths are as follows: For the root user, the path is `/usr/local/Ascend`. For non-root users, the path is `${HOME}/Ascend`, where `${HOME}` indicates the current user's directory.
	75	+The preceding environment variable configurations take effect only in the current window. You can add the `source ${HOME}/Ascend/ascend-toolkit/set_env.sh` command to the environment variable configuration file (such as the .bashrc file) as required.
	76	+
	77	+
	78	+### Installing torch_npu
	79	+
	80	+The current torch_npu version is 2.7.1.
	81	+
	82	+```bash
	83	+pip install torch_npu==2.7.1
	84	+```
	85	+
	86	+Note: If `ERROR: No matching distribution found for torch==2.7.1+cpu` is displayed, you can manually install Torch and then install torch_npu.
	87	+```bash
	88	+pip install torch==2.7.1+cpu --index-url https://download.pytorch.org/whl/cpu
	89	+```
	90	+
	91	+## Installing Triton-Ascend Using Pip
	92	+
	93	+### Latest Stable Version
	94	+You can install the latest stable version of Triton-Ascend using pip.
	95	+
	96	+```shell
	97	+pip install triton-ascend
	98	+```
	99	+
	100	+- Note: If the community edition of Triton has been installed, uninstall it first. Then install Triton-Ascend. Doing so helps prevent conflicts.
	101	+```shell
	102	+pip uninstall triton
	103	+pip install triton-ascend
	104	+```
	105	+
	106	+### Nightly Build Version
	107	+We provide daily updated nightly packages. You can run the following command to install them:
	108	+
	109	+```shell
	110	+pip install -i https://test.pypi.org/simple/ "triton-ascend<3.2.0rc" --pre --no-cache-dir
	111	+```
	112	+You can also find all nightly build packages in [History](https://test.pypi.org/project/triton-ascend/#history).
	113	+
	114	+Note: If you encounter SSL-related errors when running the `pip install` command, add the `--trusted-host test.pypi.org --trusted-host test-files.pythonhosted.org` option to solve them.
	115	+
	116	+## Installing Triton-Ascend Using the Source Code
	117	+
	118	+If you need to develop or customize Triton-Ascend, you should install it by compiling the source code. This method allows you to adjust the source code based on project requirements and compile and install a customized Triton-Ascend version.
	119	+
	120	+### System Requirements
	121	+
	122	+- GCC >= 9.4.0
	123	+- GLIBC >= 2.27
	124	+
	125	+### Dependencies
	126	+
	127	+#### Installing System Library Dependencies
	128	+
	129	+Install zlib1g-dev, LLD and Clang. You can also install ccache to accelerate the build process.
	130	+
	131	+- Recommended version: Clang >= 15
	132	+- Recommended version: LLD >= 15
	133	+
	134	+```bash
	135	+Taking Ubuntu as an example:
	136	+sudo apt update
	137	+sudo apt install zlib1g-dev clang-15 lld-15
	138	+sudo apt install ccache # optional
	139	+```
	140	+
	141	+Triton-Ascend depends heavily on zlib1g-dev. If you use the yum source, run the following installation command:
	142	+
	143	+```bash
	144	+sudo yum install -y zlib-devel
	145	+```
	146	+
	147	+#### Installing Python Dependencies
	148	+
	149	+```bash
	150	+pip install ninja cmake wheel pybind11 # build-time dependencies
	151	+```
	152	+
	153	+### Building with LLVM
	154	+
	155	+Triton uses LLVM 20 to generate code for GPUs and CPUs. Similarly, the BiSheng Compiler of Ascend depends on LLVM to generate NPU code. Therefore, you need to compile the LLVM source code. Pay attention to the specific LLVM version of dependencies. LLVM build supports two methods. You only need to follow either method.
	156	+
	157	+#### Code preparation: Run the `git checkout` command to check out the specified LLVM version.
	158	+
	159	+ ```bash
	160	+ git clone --no-checkout https://github.com/llvm/llvm-project.git
	161	+ cd llvm-project
	162	+ git checkout b5cc222d7429fe6f18c787f633d5262fac2e676f
	163	+ ```
	164	+
	165	+#### Method 1: Installing LLVM Using Clang
	166	+
	167	+- Step 1: You are advised to use Clang to install LLVM. Install Clang and LLD in the environment and specify their versions (Clang >= 15 and LLD >= 15 are recommended).
	168	+ If Clang, LLD, and ccache are not installed, run the following commands to install them:
	169	+
	170	+ ```bash
	171	+ apt-get install -y clang-15 lld-15 ccache
	172	+ ```
	173	+
	174	+- Step 2: Set the environment variable LLVM_INSTALL_PREFIX to your target installation path.
	175	+
	176	+ ```bash
	177	+ export LLVM_INSTALL_PREFIX=/path/to/llvm-install
	178	+ ```
	179	+
	180	+- Step 3: Run the following commands to build and install LLVM:
	181	+
	182	+ ```bash
	183	+ cd $HOME/llvm-project # Path to the LLVM code pulled by git clone
	184	+ mkdir build
	185	+ cd build
	186	+ cmake ../llvm \
	187	+ -G Ninja \
	188	+ -DCMAKE_C_COMPILER=/usr/bin/clang-15 \
	189	+ -DCMAKE_CXX_COMPILER=/usr/bin/clang++-15 \
	190	+ -DCMAKE_LINKER=/usr/bin/lld-15 \
	191	+ -DCMAKE_BUILD_TYPE=Release \
	192	+ -DLLVM_ENABLE_ASSERTIONS=ON \
	193	+ -DLLVM_ENABLE_PROJECTS="mlir;llvm;lld" \
	194	+ -DLLVM_TARGETS_TO_BUILD="host;NVPTX;AMDGPU" \
	195	+ -DLLVM_ENABLE_LLD=ON \
	196	+ -DCMAKE_INSTALL_PREFIX=${LLVM_INSTALL_PREFIX}
	197	+ ninja install
	198	+ ```
	199	+
	200	+#### Method 2: Installing LLVM Using GCC
	201	+
	202	+- Step 1: Clang is recommended. However, when only GCC is available, pay attention to [Note 1](#note1) and [Note 2](#note2). Set the environment variable LLVM_INSTALL_PREFIX to your target installation path.
	203	+
	204	+ ```bash
	205	+ export LLVM_INSTALL_PREFIX=/path/to/llvm-install
	206	+ ```
	207	+
	208	+- Step 2: Run the following commands to build and install LLVM:
	209	+
	210	+ ```bash
	211	+ cd $HOME/llvm-project # your clone of LLVM.
	212	+ mkdir build
	213	+ cd build
	214	+ cmake -G Ninja ../llvm \
	215	+ -DLLVM_CCACHE_BUILD=OFF \
	216	+ -DCMAKE_BUILD_TYPE=Release \
	217	+ -DLLVM_ENABLE_ASSERTIONS=ON \
	218	+ -DLLVM_ENABLE_PROJECTS="mlir;llvm" \
	219	+ -DLLVM_TARGETS_TO_BUILD="host;NVPTX;AMDGPU" \
	220	+ -DCMAKE_INSTALL_PREFIX=${LLVM_INSTALL_PREFIX}
	221	+ ninja install
	222	+ ```
	223	+
	224	+<a id="note1"></a>Note 1: If `ld.lld: error: undefined symbol` is displayed during compilation, you can add `-DLLVM_ENABLE_LLD=ON` during Step 2.
	225	+
	226	+<a id="note2"></a>Note 2: If ccache has been installed and is running properly in the environment, you can set `-DLLVM_CCACHE_BUILD=ON` to accelerate the build process. Otherwise, do not enable it.
	227	+
	228	+#### Cloning Triton-Ascend
	229	+
	230	+```bash
	231	+git clone https://gitcode.com/Ascend/triton-ascend.git && cd triton-ascend/python
	232	+```
	233	+
	234	+#### Building Triton-Ascend
	235	+
	236	+1. Install the source code.
	237	+
	238	+- Step 1: Ensure that the target installation path of LLVM (${LLVM_INSTALL_PREFIX}) has been set in the [Building with LLVM] section.
	239	+- Step 2: Ensure that Clang 15 or later, LLD 15 or later, and ccache have been installed.
	240	+
	241	+ ```bash
	242	+ LLVM_SYSPATH=${LLVM_INSTALL_PREFIX} \
	243	+ TRITON_BUILD_WITH_CCACHE=true \
	244	+ TRITON_BUILD_WITH_CLANG_LLD=true \
	245	+ TRITON_BUILD_PROTON=OFF \
	246	+ TRITON_WHEEL_NAME="triton-ascend" \
	247	+ TRITON_APPEND_CMAKE_ARGS="-DTRITON_BUILD_UT=OFF" \
	248	+ python3 setup.py install
	249	+ ```
	250	+
	251	+- Note 3: GCC 9.4.0 or later is recommended. If the GCC version is earlier than 9.4.0, "ld.lld: error: unable to find library -lstdc++fs" may be reported, indicating that the linker cannot find the stdc++fs library.
	252	+This library supports the file system features of versions earlier than GCC 9. In this case, you need to manually uncomment the related code snippet in the CMake file.
	253	+
	254	+- triton-ascend/CMakeLists.txt
	255	+
	256	+ ```bash
	257	+ if (NOT WIN32 AND NOT APPLE)
	258	+ link_libraries(stdc++fs)
	259	+ endif()
	260	+ ```
	261	+
	262	+ After uncommenting the code snippet, rebuild the project to solve the problem.
	263	+
	264	+2. Run the Triton example.
	265	+
	266	+ Install the runtime dependencies. Refer to the following command:
	267	+ ```bash
	268	+ cd triton-ascend && pip install -r requirements_dev.txt
	269	+ ```
	270	+ Run the [01-vector-add.py](../../ascend/examples/tutorials/01-vector-add.py) instance.
	271	+ ```bash
	272	+ # Set the CANN environment variables (for example, as the root user and with the default installation path /usr/local/Ascend).
	273	+ source /usr/local/Ascend/ascend-toolkit/set_env.sh
	274	+ # Run the tutorials example.
	275	+ python3 ./triton-ascend/ascend/examples/tutorials/01-vector-add.py
	276	+ ```
	277	+ If an output similar to the following is displayed, the environment is correctly configured:
	278	+ ```
	279	+ tensor([0.8329, 1.0024, 1.3639, ..., 1.0796, 1.0406, 1.5811], device='npu:0')
	280	+ tensor([0.8329, 1.0024, 1.3639, ..., 1.0796, 1.0406, 1.5811], device='npu:0')
	281	+ The maximum difference between torch and triton is 0.0
	282	+ ```

Adocs/en/migration_guide/architecture_difference.md+159-0

@@ -0,0 +1,159 @@
	1	+# Development Differences Between Ascend and GPUs
	2	+
	3	+## Multi-Core Task Parallelism Strategy
	4	+
	5	+NPUs are strongly bound to physical cores in Triton multi-core parallelism. This represents a core difference from GPUs' logical dimension parallelism + automatic physical mapping in hardware.
	6	+
	7	+- Core comparison
	8	+
	9	+ \|Dimension \|GPU (NVIDIA)\|Ascend\|
	10	+ \|-----------\|--------------\|-----------\|
	11	+ \|Essence of grids\| Logical task dimension (decoupled from physical cores)\| Physical core group mapping (bound to the AI core topology)\|
	12	+ \|Limit on the number of cores/dimensions\| No hard limit on the grid dimensions/sizes\| Grid size ≤ Total number of AI cores; topology matching required by 2D\|
	13	+
	14	+GPUs can be bound to multiple dimensions (a 3D grid of `[n, m, l]` is equivalent to` n × m × l` parallel threads). Each thread corresponds to only one kernel execution and executes only once.\
	15	+In NPUs, vector cores and cube cores belong to multiple physical cores. The number of cores varies with the generation of hardware. Each core executes only one block and can schedule the block execution repeatedly.
	16	+
	17	+### Full Utilization of Cores
	18	+
	19	+Ascend NPUs have multiple computing cores. Properly allocating and fully utilizing all available cores is one of the key factors to improve operator performance.
	20	+When calling Triton kernel functions, you can set the launch parameter to control the number of cores in use. Take the GELU operator as an example:
	21	+
	22	+```Python
	23	+triton_gelu[n, 1, 1](...) # The first parameter indicates the number of cores in use. n indicates that n cores are in use.
	24	+```
	25	+
	26	+By optimizing the number of cores, you can fully schedule and utilize all computing resources, thereby maximizing the degree of parallelism (DOP) and throughput. Note that the number of cores in the current version must be less than or equal to 65,535.
	27	+
	28	+
	29	+## Single-Core Data Transfer Strategy
	30	+
	31	+### Data Tiling
	32	+
	33	+When you write Triton kernel functions, a proper data tiling strategy is essential for performance optimization. By adjusting tiling granularity parameters, you can balance computational workload and memory access efficiency across different dimensions.
	34	+
	35	+Common tiling parameters include:
	36	+
	37	+```text
	38	+ncore: the number of cores in use (cross-core tiling)
	39	+xblock: the size of inter-core data blocks (inter-core tiling)
	40	+xblock_sub: the granularity of intra-core tiling (fine-grained intra-core tiling)
	41	+```
	42	+
	43	+By manually selecting the optimal tiling configurations based on your actual scenario, you can maximize the utilization of on-chip memory during each computation cycle, preventing performance bottlenecks caused by frequent access to the global memory.
	44	+
	45	+
	46	+Taking the GELU operator as an example, adjusting the tiling parameters helps effectively adapt to the on-chip cache capacity limit, thereby improving execution efficiency.
	47	+
	48	+Note: Atlas 800T/I A2 has an on-chip memory capacity of 192 KB. When designing the tiling strategy, ensure that the data volume of each computation cycle does not exceed this capacity.
	49	+
	50	+
	51	+#### Example GELU Operator
	52	+
	53	+The following demonstrates the development of an example GELU operator with three result computation methods.
	54	+
	55	+`standard_unary` is standard Torch computation.
	56	+
	57	+`triton_easy_kernel` is a simple implementation of Triton.
	58	+
	59	+`triton_better_kernel` is a more efficient implementation of Triton.
	60	+
	61	+#### Standard Torch Writing
	62	+
	63	+After computing the input `tensor x0`, Torch implements the GELU operator and returns the result value.
	64	+
	65	+```Python
	66	+def standard_unary(x0):
	67	+ res = x0 * 0.5 * (1.0 + torch.erf(x0 / torch.sqrt(torch.tensor(2.0))))
	68	+ return res
	69	+```
	70	+
	71	+#### Simple Triton Writing
	72	+
	73	+The following is an example of a simple kernel written in Triton, demonstrating how to define and call a basic Triton kernel function. This example implements a simple mathematical operation (GELU activation function).
	74	+
	75	+```Python
	76	+# Define the triton_kernel function.
	77	+@triton.jit
	78	+def triton_easy_kernel(in_ptr0, out_ptr0, NUMEL: tl.constexpr):
	79	+ idx_block = tl.arange(0, NUMEL)
	80	+ x = tl.load(in_ptr0 + idx_block)
	81	+ ret = x * 0.5 * (1.0 + tl.erf(x / tl.sqrt(2.0)))
	82	+ tl.store(out_ptr0 + idx_block, ret)
	83	+```
	84	+
	85	+Precautions
	86	+
	87	+1. Memory limit: In the preceding writing, all input data is loaded to memory at a time for computation. If the input tensor is too large, it may exceed the on-chip memory capacity of a single kernel, resulting in a memory overflow error.
	88	+Therefore, this simple writing is suitable for computing small-scale tensors or for understanding the basic writing and call method of Triton kernels.
	89	+
	90	+2. Application scenarios: This method helps developers quickly understand and get started with Triton programming. However, for large-scale data sets or scenarios demanding high performance, developers are advised to use more complex data tiling strategies to fully utilize hardware resources and prevent memory overflow. In this way, developers can quickly get started with Triton programming and understand how to define, call, and optimize Triton kernel functions.
	91	+
	92	+
	93	+#### More Efficient Triton Writing
	94	+
	95	+When using Triton to write high-performance operators on Ascend NPUs, developers need to use a data tiling strategy to fully utilize hardware resources, prevent memory overflow, and improve execution efficiency.
	96	+The following is an example of an optimized Triton kernel implementation suitable for large-scale tensor computation.
	97	+
	98	+```Python
	99	+# Define the triton_kernel function.
	100	+@triton.jit
	101	+def triton_better_kernel(in_ptr0, out_ptr0, xnumel, XBLOCK: tl.constexpr, XBLOCK_SUB: tl.constexpr):
	102	+ xoffset = tl.program_id(0) * XBLOCK
	103	+ for xoffset_sub in range(0, XBLOCK, XBLOCK_SUB):
	104	+ x_index = xoffset + xoffset_sub + tl.arange(0, XBLOCK_SUB)[:]
	105	+ xmask = x_index < xnumel
	106	+ x = tl.load(in_ptr0 + x_index, xmask)
	107	+ ret = x * 0.5 * (1.0 + tl.erf(x / tl.sqrt(2.0)))
	108	+ tl.store(out_ptr0 + x_index, ret, xmask)
	109	+
	110	+# Call the triton_kernel function.
	111	+ncore = 32
	112	+xblock = 32768
	113	+xblock_sub = 8192
	114	+triton_better_kernel[ncore, 1, 1](x0, out1, x0.numel(), xblock, xblock_sub)
	115	+```
	116	+
	117	+Explanation of key code:
	118	+
	119	+```Python
	120	+# Calculate the start offset address of the data block processed by the current core to implement inter-core tiling. Each core is responsible only for a data segment of size XBLOCK.
	121	+xoffset = tl.program_id(0) * XBLOCK
	122	+
	123	+# Further split the data block within a single core to process data of size XBLOCK_SUB each time, which is known as intra-core tiling.
	124	+for xoffset_sub in range(0, XBLOCK, XBLOCK_SUB):
	125	+
	126	+# Construct the data index array of the current iteration. This array is used to access the input and output tensors.
	127	+x_index = xoffset + xoffset_sub + tl.arange(0, XBLOCK_SUB)[:]
	128	+
	129	+# Set a mask to prevent out-of-bounds access and ensure that only data within the defined range is processed.
	130	+xmask = x_index < xnumel
	131	+
	132	+# Load data from the global memory to the on-chip memory and write the computation results back to the global memory.
	133	+tl.load() and tl.store()
	134	+```
	135	+
	136	+## Compilation Optimization
	137	+### Ascend NPU IR Optimization
	138	+The following table lists the compilation options for Ascend NPU IR optimization, which are adapted to the hardware and software features of Ascend.
	139	+Usage: During the autotune configuration phase, pass the values of the compilation options.
	140	+For example, to enable the `multibuffer` option, pass `'multibuffer': True` to `triton.Config` during the autotune configuration phase. For details, see [Autotune Example](../examples/06_autotune_example.md).
	141	+```python
	142	+ def get_autotune_config():
	143	+ return [
	144	+ triton.Config({'XS': 1 * 128, 'multibuffer': True}),]
	145	+```
	146	+
	147	+\| Option \| Capability \| Enabled or Not\|
	148	+\| ----------------- \| ------------ \| ----------------- \|
	149	+\| multibuffer \| Data transfer through parallel pipelines. \| Default: true. Options: true and false. It is configurable during autotune. \|
	150	+\| unit_flag \| Optimization item for cube-out. \| Default: None. Options: true and false. It is configurable during autotune. \|
	151	+\| limit_auto_multi_buffer_only_for_local_buffer \| Optimization item for CV operators and cube-out. \| Default: None. Options: true and false. It is configurable during autotune.\|
	152	+\| limit_auto_multi_buffer_of_local_buffer \| Scope of enabling double buffer for cube operators. \| Default: None. Value range: ["no-limit","no-l0c"]. It is configurable during autotune. \|
	153	+\| set_workspace_multibuffer \| It takes effect only when limit_auto_multi_buffer_only_for_local_buffer is set to false.\| Default: None. Example: [2,4]. It is configurable during autotune. \|
	154	+\| enable_hivm_auto_cv_balance \| set_workspace_multibuffer takes effect only when limit_auto_multi_buffer_only_for_local_buffer is set to false.\| Default: None. Options: true and false. It is configurable during autotune.\|
	155	+\| tile_mix_vector_loop \| Optimization item for CV operators. It specifies the number of segments into which the current vector can be split. \| Default: None. Example: [2,4,8]. It is configurable during autotune. \|
	156	+\| tile_mix_cube_loop \| Optimization item for CV operators. It specifies the number of segments into which the current cube can be split. \| Default: None. Example: [2,4,8]. It is configurable during autotune. \|
	157	+
	158	+- Note: The compilation optimization options are located in ascend/backend/compiler.py.
	159	+- Note: CV operators indicate that both AI cores and vector cores are used during operator computation.

Adocs/en/migration_guide/migrate_from_gpu.md+339-0

@@ -0,0 +1,339 @@
	1	+# Migrating Triton Operators from GPUs
	2	+This document outlines key considerations for migrating Triton operators from GPUs, organized into three key aspects: multi-core task parallelism, single-core data transfer, and single-core data computation. In the "Multi-Core Task Parallelism" section, we highlight core migration principles and provide a complete migration example. The "Single-Core Data Transfer" section describes the basic procedure for migrating operators from GPUs to NPUs. Finally, the "Single-Core Data Computation" analyzes the differences between GPUs and NPUs with respect to Triton issues. We also provide answers to frequently asked questions (FAQ).
	3	+
	4	+## Multi-Core Task Parallelism
	5	+
	6	+### Core Migration Principles
	7	+- Shift from GPUs' "logical grid flexibility" to Ascend's "physical core group binding".
	8	+- Enforce 32-byte memory alignment in VV scenarios and 512-byte memory alignment in CV scenarios. Remove GPU-specific synchronization APIs.
	9	+- Prefer 1D grids. NPUs' 2D adaptations will be merged into the 1D form. Actual grid values must align with the physical core count available on chips. For example, `(20,)` and `(4, 5)` will produce equivalent execution results.
	10	+
	11	+### Complete Migration Example (Vector Addition)
	12	+
	13	+```diff
	14	++ import torch_npu # [Added] Import Ascend NPUs' PyTorch adaptation library to support NPU devices.
	15	+import triton
	16	+import triton.language as tl
	17	+
	18	+- DEVICE = triton.runtime.driver.active.get_active_torch_device() # [Deleted] GPU devices are automatically obtained. NPUs do not need this logic.
	19	+
	20	+@triton.jit
	21	+def add_kernel(x_ptr, # Pointer to first input vector.
	22	+y_ptr, # Pointer to second input vector.
	23	+output_ptr, # Pointer to output vector.
	24	+n_elements, # Size of the vector.
	25	+BLOCK_SIZE: tl.constexpr, # Number of elements each program should process.
	26	+):
	27	+ pid = tl.program_id(axis=0) # We use a 1D launch grid so axis is 0.
	28	+ block_start = pid * BLOCK_SIZE
	29	+ offsets = block_start + tl.arange(0, BLOCK_SIZE)
	30	+ mask = offsets < n_elements
	31	+ x = tl.load(x_ptr + offsets, mask=mask)
	32	+ y = tl.load(y_ptr + offsets, mask=mask)
	33	+ output = x + y
	34	+ tl.store(output_ptr + offsets, output, mask=mask)
	35	+
	36	+def add(x: torch.Tensor, y: torch.Tensor):
	37	+ output = torch.empty_like(x)
	38	+- assert x.device == DEVICE and y.device == DEVICE and output.device == DEVICE # [Deleted] GPU devices have consistency checks. NPUs do not need explicit assertion.
	39	+ n_elements = output.numel()
	40	+ grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']), )
	41	+ add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=1024)
	42	+ return output
	43	+
	44	+torch.manual_seed(0)
	45	+size = 98432
	46	+- x = torch.rand(size, device='cuda') # [Deleted] Specify the GPU device.
	47	++ x = torch.rand(size, device='npu') # [Modified] Specify the Ascend NPU device.
	48	+- y = torch.rand(size, device='cuda') # [Deleted] Specify the GPU device.
	49	++ y = torch.rand(size, device='npu') # [Modified] Specify the Ascend NPU device.
	50	+output_torch = x + y
	51	+output_triton = add(x, y)
	52	+print(output_torch)
	53	+print(output_triton)
	54	+print(f'The maximum difference between torch and triton is '
	55	+f'{torch.max(torch.abs(output_torch - output_triton))}')
	56	+```
	57	+
	58	+
	59	+## Single-Core Data Transfer
	60	+First off, you need to understand the basic steps for migrating from GPUs to NPUs. Below is an example Triton kernel that runs properly on GPUs:
	61	+- The first step is to change `device='cuda'` to `device='npu'` to run the kernel on NPUs.
	62	+
	63	+```diff
	64	+import pytest
	65	+import torch
	66	+import triton
	67	+import triton.language as tl
	68	+
	69	+@triton.jit
	70	+def fn_broadcast_1d(output_ptr, x_ptr, XS: tl.constexpr, YS: tl.constexpr):
	71	+ xidx = tl.arange(0, XS)[None, :]
	72	+ base = tl.load(x_ptr + xidx)
	73	+ out = base.broadcast_to((YS, XS))
	74	+ oidx = tl.arange(0, YS)[:, None] * XS + tl.arange(0, XS)[None, :]
	75	+ tl.store(output_ptr + oidx, out)
	76	+
	77	+@pytest.mark.parametrize('shape', [(1,), (2,), (4,)])
	78	+@pytest.mark.parametrize('dtype', [torch.int32])
	79	+def test_npu_1d(shape, dtype):
	80	+ XS = shape[0]
	81	+ YS = 4
	82	+
	83	+- x = torch.randint(-1000, 1000, (XS,), dtype=dtype, device='cuda')
	84	++ x = torch.randint(-1000, 1000, (XS,), dtype=dtype, device='npu')
	85	+ std = torch.broadcast_to(x, (YS, XS))
	86	+- output = torch.randint(-1000, 1000, (YS, XS), dtype=dtype, device='cuda')
	87	++ output = torch.randint(-1000, 1000, (YS, XS), dtype=dtype, device='npu')
	88	+ fn_broadcast_1d[(1,)](output, x, XS, YS)
	89	+ assert torch.allclose(std, output)
	90	+```
	91	+
	92	+
	93	+## Single-Core Data Computation
	94	+
	95	+
	96	+### Difference Analysis
	97	+Ascend NPUs are equipped with multiple compute cores (AI cores), categorized into cube cores and vector cores. The exact number of AI cores varies by chip model and can be queried through the driver.active.utils.get_device_properties API. When executing a Triton kernel, the runtime APIs allow the number of concurrent tasks to exceed the available physical AI cores—though the total number of concurrent tasks is capped at 65,535. In such cases, these tasks are divided into multiple batches and scheduled to NPUs for execution. Crucially, the number of concurrent tasks within each individual batch still cannot surpass the number of physical AI cores. This batch scheduling introduces additional device-side overhead, which can impact the overall execution performance of Triton operators.
	98	+
	99	+To maximize the utilization of physical AI core resources on NPUs for accelerated parallel computing and minimize batch scheduling overhead, it is advisable to set the number of concurrent tasks to match the number of the underlying physical AI cores. For Triton operators that perform only vector core computations, the number of concurrent tasks should be equal to the number of vector cores. For other types of Triton operators (those using tl.dot), the number of concurrent tasks should be equal to the total number of AI cores.
	100	+Tips: TRITON_ALL_BLOCKS_PARALLEL controls the automatic optimization of the number of logical cores based on the number of physical cores. This feature can be enabled only when logical cores can execute in parallel. When the number of logical cores is greater than the number of physical cores, enabling this feature will instruct the compiler to automatically adjust the number of logical cores to match the number of physical cores, thereby reducing scheduling overhead.
	101	+\|Dimension \| Core Structure \| Operator Type \|
	102	+\|--------\|----------------------------------------\|------------------------------------------------\|
	103	+\|Ascend NPU\|Multiple AI cores, categorized into cube cores (for matrix multiplication) and vector cores (for vector computation)\| Vector-only operators → Number of concurrent tasks = Number of vector cores; Operators using tl.dot → Number of concurrent tasks = Number of AI cores\|
	104	+\|GPU NVIDIA/AMD\| Multiple CUDA cores (for scalar/vector computation) + Tensor cores (for matrix multiplication)\| Generally, GPU operators can be mapped to CUDA cores or tensor cores. The concurrency is automatically determined by the compiler and hardware.\|
	105	+
	106	+
	107	+## FAQ
	108	+After completing the basic migration procedure, you may encounter the following two types of new issues:
	109	+1. coreDim limit
	110	+This issue is triggered when grid dimensions exceed the hardware limit of NPUs.
	111	+Typical error message: `coreDim=xxxx can't be greater than UINT16_MAX`.
	112	+2. UB space overflow
	113	+Memory usage exceeds the NPU cache capacity.
	114	+Typical error message: `ub overflow, requires xxxx bits while 1572684 bits available!`.
	115	+
	116	+
	117	+### Solving the coreDim Limit Issue
	118	+Issue analysis:
	119	+The coreDim parameter of NPUs cannot exceed UINT16_MAX (65535). When processing large-scale data, simplistic grid division may exceed this limit.
	120	+
	121	+Case: Optimizing the `zeros_like` function
	122	+(data scale `N = 1073741824`; original `BLOCK_SIZE = 2048`; calculated `coreDim = 524288`, exceeding the limit of 65535)
	123	+
	124	+Solution 1:
	125	+To address the coreDim limit in the Ascend compiler, one solution is to set the environment variable 'TRITON_ALL_BLOCKS_PARALLEL' to 1 by running this command:
	126	+export TRITON_ALL_BLOCKS_PARALLEL=1
	127	+Solution 2:
	128	+Another solution is to increase BLOCK_SIZE to reduce the number of required cores and ensure that coreDim remains within the limit.
	129	+The calculation follows: `coreDim = ceil(N / BLOCK_SIZE)`. → It needs to satisfy `ceil(N / BLOCK_SIZE) <= 65535 => BLOCK_SIZE >= ceil(N / 65535)`. Given `N = 1073741824`, we have `BLOCK_SIZE >= triton.next_power_of_2(triton.cdiv(1073741824, 65535)) = 32768`. Therefore, 32768 is the minimum safe value.
	130	+
	131	+Code before optimization:
	132	+```diff
	133	+import logging
	134	+import torch
	135	+import triton
	136	+import triton.language as tl
	137	+logger = logging.getLogger(name)
	138	+@triton.jit
	139	+def zeros_kernel(
	140	+ output_ptr,
	141	+ n_elements,
	142	+ BLOCK_SIZE: tl.constexpr,
	143	+ ):
	144	+ pid = tl.program_id(axis=0)
	145	+ block_start = pid * BLOCK_SIZE
	146	+ offsets = block_start + tl.arange(0, BLOCK_SIZE)
	147	+ mask = offsets < n_elements
	148	+ tl.store(output_ptr + offsets, 0.0, mask=mask)
	149	+
	150	+def zeros_like(x, *, dtype=None, layout=None, device=None, pin_memory=None, memory_format=None):
	151	+ logger.debug("GEMS ZEROS_LIKE")
	152	+ if device is None:
	153	+ device = x.device # x.device = "npu"
	154	+ if dtype is None:
	155	+ dtype = x.dtype
	156	+
	157	+ out = torch.empty_like(x, device=device, dtype=dtype)
	158	+ N = x.numel()
	159	+ grid_fn = lambda meta: (triton.cdiv(N, meta["BLOCK_SIZE"]),)
	160	+
	161	+ zeros_kernel[grid_fn](out, N, BLOCK_SIZE=1024) # The original value is too small.
	162	+ return out
	163	+```
	164	+Code after optimization:
	165	+```diff
	166	+import logging
	167	+import torch
	168	+import triton
	169	+import triton.language as tl
	170	+logger = logging.getLogger(name)
	171	+@triton.jit
	172	+def zeros_kernel(
	173	+ output_ptr,
	174	+ n_elements,
	175	+ BLOCK_SIZE: tl.constexpr,
	176	+ ):
	177	+ pid = tl.program_id(axis=0)
	178	+ block_start = pid * BLOCK_SIZE
	179	+ offsets = block_start + tl.arange(0, BLOCK_SIZE)
	180	+ mask = offsets < n_elements
	181	+ tl.store(output_ptr + offsets, 0.0, mask=mask)
	182	+
	183	+def zeros_like(x, *, dtype=None, layout=None, device=None, pin_memory=None, memory_format=None):
	184	+ logger.debug("GEMS ZEROS_LIKE")
	185	+ if device is None:
	186	+ device = x.device # x.device = "npu"
	187	+ if dtype is None:
	188	+ dtype = x.dtype
	189	+
	190	+ out = torch.empty_like(x, device=device, dtype=dtype)
	191	+ N = x.numel()
	192	+ min_block_size = triton.next_power_of_2(triton.cdiv(N, 65535))
	193	+ BLOCK_SIZE = max(32768, min_block_size) # The minimum value is 32768.
	194	+ grid_fn = lambda meta: (triton.cdiv(N, meta["BLOCK_SIZE"]),)
	195	+
	196	+ zeros_kernel[grid_fn](out, N, BLOCK_SIZE=BLOCK_SIZE)
	197	+ return out
	198	+```
	199	+
	200	+### Dynamically Calculating BLOCK_SIZE to Ensure coreDim Remains Within the Limit
	201	+```diff
	202	+optimal_block_size = 32768 # Optimized value obtained after calculation
	203	+
	204	+grid_fn = lambda meta: (triton.cdiv(N, optimal_block_size),)
	205	+
	206	+zeros_kernel[grid_fn](out, N, BLOCK_SIZE=optimal_block_size)
	207	+return out
	208	+```
	209	+
	210	+### Handling the Compound Issue: coreDim + UB Overflow
	211	+Issue analysis:
	212	+In some scenarios, solving the coreDim limit issue may inadvertently trigger a new issue—UB overflow. This typically occurs when increasing BLOCK_SIZE causes the data volume processed by a single thread block to exceed the UB cache capacity of NPUs.
	213	+
	214	+Case:
	215	+Data scale `N = 1073741824`; original `BLOCK_SIZE = 4096`; calculated `coreDim = 262144`, exceeding the limit of 65535. After BLOCK_SIZE is adjusted to 32768, coreDim is 32768 (within the limit), but UB overflow occurs.
	216	+
	217	+Solution:
	218	+Introduce the BLOCK_SIZE_SUB parameter to further subdivide large blocks, thereby controlling memory usage while maintaining a reasonable coreDim.
	219	+Code before optimization:
	220	+```diff
	221	+import logging
	222	+import torch
	223	+import triton
	224	+import triton.language as tl
	225	+logger = logging.getLogger(name)
	226	+
	227	+@triton.jit
	228	+def masked_fill_kernel(inp, expand_mask, value, out, N, BLOCK_SIZE: tl.constexpr):
	229	+ pid = tl.program_id(axis=0)
	230	+ offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
	231	+ mask = offsets < N
	232	+ fill_mask = tl.load(expand_mask + offsets, mask=mask, other=0).to(tl.int1)
	233	+ cur_inp = tl.load(inp + offsets, mask=(~fill_mask) & mask, other=0)
	234	+ tl.store(out + offsets, cur_inp, (~fill_mask) & mask)
	235	+ tl.store(out + offsets, value, fill_mask & mask)
	236	+ def masked_fill(inp, mask, value):
	237	+ # ... Parameter verification code ...
	238	+ # inp.device = "npu"
	239	+ N = inp.numel()
	240	+ if N == 0:
	241	+ return out
	242	+
	243	+ grid = lambda meta: (triton.cdiv(N, 4096),) # coreDim exceeds the limit.
	244	+ masked_fill_kernel[grid](inp, mask.to(torch.int), value, out, N, 4096)
	245	+ return out
	246	+```
	247	+Code after optimization:
	248	+```diff
	249	+import logging
	250	+import torch
	251	+import triton
	252	+import triton.language as tl
	253	+logger = logging.getLogger(name)
	254	+
	255	+@triton.jit
	256	+def masked_fill_kernel(inp, expand_mask, value, out, N,
	257	+ BLOCK_SIZE: tl.constexpr, BLOCK_SIZE_SUB: tl.constexpr):
	258	+ pid = tl.program_id(axis=0)
	259	+ base_offset = pid * BLOCK_SIZE
	260	+ # Calculate the number of sub-blocks to be processed.
	261	+ num_sub_blocks = tl.cdiv(BLOCK_SIZE, BLOCK_SIZE_SUB)
	262	+ # Process blocks to avoid UB overflow.
	263	+ for sub_block_idx in range(num_sub_blocks):
	264	+ sub_offset = base_offset + sub_block_idx * BLOCK_SIZE_SUB
	265	+ offsets = sub_offset + tl.arange(0, BLOCK_SIZE_SUB)
	266	+ mask = offsets < N
	267	+ # Load and process data in batches.
	268	+ input_vals = tl.load(inp + offsets, mask=mask, other=0)
	269	+ fill_mask_vals = tl.load(expand_mask + offsets, mask=mask, other=0).to(tl.int1)
	270	+ # First, write the original data.
	271	+ tl.store(out + offsets, input_vals, mask=mask)
	272	+ # Then overwrite the target value at the position where padding is required.
	273	+ value_to_write = tl.full([BLOCK_SIZE_SUB], value, dtype=input_vals.dtype)
	274	+ final_vals = tl.where(fill_mask_vals, value_to_write, input_vals)
	275	+ tl.store(out + offsets, final_vals, mask=mask)
	276	+
	277	+def masked_fill(inp, mask, value):
	278	+ logger.debug("GEMS MASKED FILL")
	279	+
	280	+ # ... Parameter verification code ...
	281	+ # inp.device = "npu"
	282	+ N = inp.numel()
	283	+ if N == 0:
	284	+ return out
	285	+
	286	+ # Use optimized parameter settings.
	287	+ MAIN_BLOCK_SIZE = 32768 # Ensure that coreDim is within the limit.
	288	+ SUB_BLOCK_SIZE = 1024 # Control the UB usage.
	289	+
	290	+ grid = lambda meta: (triton.cdiv(N, MAIN_BLOCK_SIZE),)
	291	+ masked_fill_kernel[grid](inp, expand_mask.to(torch.int), value, out, N,
	292	+ MAIN_BLOCK_SIZE, SUB_BLOCK_SIZE)
	293	+ return out
	294	+```
	295	+
	296	+### Why Does the UBSIZE Out of Memory Error Occur?
	297	+Improper data tiling can lead to excessive unaligned memory access or computation. Consider a 2D data transfer of shape `(64, 32)` as an example. The corresponding stride is `(12832, 128)`. If aligned memory access is required, the stride becomes `(32, 1)`. In unaligned access scenarios, an additional axis of size `1` is added to the innermost dimension, yielding a shape of `(64, 32, 4)`. Because the hardware mandates 32-byte UB memory alignment in VV scenarios, the corresponding stride is recalculated as `(12832, 128, 1)`, assuming `type=float16`.
	298	+
	299	+
	300	+### Discrete Memory Access and Inefficient Scalar Mapping Observed by Line-by-Line Code Comparison
	301	+Set the environment variable TRITON_DEBUG to 1, save ~/.triton/cache/xxx.ttadapter, and execute:
	302	+```diff
	303	+bishengir-compile xxx.ttadapter --target=Ascend910B3 --enable-auto-multi-buffer=True --enable-hfusion-compile=true --enable-hivm-compile=true --enable-triton-kernel-compile=true --hivm-compile-args=bishengir-print-ir-after=hivm-inject-sync
	304	+```
	305	+Compare the Triton-Python algorithm's logic with the internal operations of the output intermediate representations (IRs) to identify any operations that are not mapped to instructions.
	306	+Check whether pure scalar transfer or computation exists in the HIVM IR phase without being mapped to SIMD instructions. If such cases exist, they will create a significant performance bottleneck.
	307	+
	308	+Problem: Discrete memory access and inefficient scalar mapping
	309	+Given `b[1024, 32] = a[1024, 32]`, the original Triton code binds thread blocks to the lowest dimension `32` in `[1024, 32]`, and then splits `1024` into `16` parts, yielding `[64, 16, 32]`. Finally, it binds thread blocks to dimension `64`.
	310	+```diff
	311	+chunk_fwd_kernel_o[(NT, B * H)](
	312	+ p_g = tl.make_block_ptr(g, (T,), (H,), (i_t * BT,), (BT,), (0,))
	313	+ block_ptr = tl.make_block_ptr(
	314	+ base=input_ptr,
	315	+ shape=(1024,), # 1D tensor
	316	+ strides=(32,), # Contiguous memory
	317	+ offsets=(i_t * 16,), # Start position
	318	+ block_shape=(BT,), # Block size
	319	+ order=(0,) # Sequential access
	320	+ )
	321	+)
	322	+```
	323	+
	324	+Optimization Approach
	325	+
	326	+Adjust shape and stride for block_ptr as follows:
	327	+The shape (1024, 32) is treated as a 2D matrix, where the lowest dimension `32` is contiguous. Accordingly, the stride should be `(32, 1)` instead of `(32,)`. This enables each thread block to access 32 contiguous elements. Bind thread blocks to the row dimension `(1024)` and configure each thread to process all 32 elements in a row. This approach guarantees contiguous memory access and high memory affinity
	328	+
	329	+Example:
	330	+```diff
	331	+block_ptr = tl.make_block_ptr(
	332	+ base=input_ptr,
	333	+ shape=(1024, 32),
	334	+ strides=(32, 1),
	335	+ offsets=(i_t * BT, 0),
	336	+ block_shape=(BT, 32),
	337	+ order=(1, 0) # First row then column (FRTC)
	338	+)
	339	+```

Adocs/en/migration_guide/performance_guidelines.md+213-0

@@ -0,0 +1,213 @@
	1	+# NPU High-Performance Programming Guide
	2	+
	3	+## Combining Grid Cores
	4	+### I. Principles for Automatically Combining Grid Cores
	5	+
	6	+Some scenarios requiring migration of Triton operators from GPUs to NPUs. Due to architectural differences, the Triton operators developed on GPUs often utilize large grid core counts. When executed on NPUs, these operators cannot be scheduled all at once. Delivering them in batches introduces significant latency and degrades performance. To optimize NPU-based Triton operators, you need to check the grid core counts first. In cases with large grid core counts, set the environment variable TRITON_ALL_BLOCKS_PARALLEL to improve operator execution performance.
	7	+
	8	+## Optimizing Instruction Parallelism
	9	+
	10	+### I. Core Principles of Instruction Parallelism Optimization
	11	+
	12	+When executing Triton operators, NPUs leverage parallel mechanisms such as multi-buffer and instruction parallelism to parallelize data-in, computation, and data-out, thereby enhancing performance. However, in certain scenarios, the multi-buffer mechanism cannot be enabled, which reduces the degree of parallelism (DOP) and degrades operator execution performance. If this issue occurs during performance optimization, consider the following aspects and implement optimizations based on the provided code examples:\
	13	+1. Data transfer and computation involve dependencies, which introduce synchronization. The memory transfer engine (MTE) can only be triggered after vector computation completes, resulting in low DOP.\
	14	+2. In cases where the operator lacks multiple data loads or a single execution completes without tiling, the multi-buffer mechanism cannot be enabled.\
	15	+3. The multi-buffer mechanism requires additional UB space. If the UB space is insufficient during computation, the multi-buffer mechanism cannot be enabled.
	16	+
	17	+### II. Code Examples
	18	+
	19	+- Example 1: Reducing synchronization for higher DOP
	20	+
	21	+ In operator optimization, increasing instruction-level parallelism (DOP) is a critical strategy. In the `tl.load` statement below, when `N` > `M`, the loaded data fills only a portion of the tensor memory space pointed to by `data`. For the remaining unfilled portion, if users do not specify the `other` value, GPUs default to zero-padding. To reduce the adaptation workload of migration, NPUs maintain the same behavior as GPUs. NPUs first use the vector core to set all the memory space pointed to by `data` to a specified value (defaulting to `0` if no `other` value is provided). Subsequently, the MTE2 instruction transfers data to part of the memory space pointed to by `data`. This implementation results in a dependency between MTE2 and vector operations, which limits parallelism and degrades overall performance.
	22	+
	23	+ ```diff
	24	+ @triton.jit
	25	+ def npu_vector_add_kernel(
	26	+ input, # [Tensor] input tensor (1 x col)
	27	+ output, # [Tensor] output tensor (1 x col)
	28	+ M: tl.constexpr, # len of the vector
	29	+ BLOCK_SIZE: tl.constexpr
	30	+ ):
	31	+ N :tl.constexpr = BLOCK_SIZE
	32	+ idx = tl.arange(0, N)
	33	+ mask = idx < M
	34	+ data = tl.load(input + idx, mask = mask) # Alternatively, specify a value such as other=-1.
	35	+ ```
	36	+
	37	+ To increase DOP and enhance performance, when the loaded data fills only a portion of the memory space pointed to by `data`, add `care_padding=False` to the load statement to remove default-value padding, provided that the unfilled portion does not affect subsequent computation results. That is, the preceding operator can be optimized follows:
	38	+
	39	+ ```diff
	40	+ @triton.jit
	41	+ def npu_vector_add_kernel(
	42	+ input, # [Tensor] input tensor (1 x col)
	43	+ output, # [Tensor] output tensor (1 x col)
	44	+ M: tl.constexpr, # len of the vector
	45	+ BLOCK_SIZE: tl.constexpr
	46	+ ):
	47	+ idx = tl.arange(0, N)
	48	+ mask = idx < M
	49	+ - data = tl.load(input + idx, mask = mask) # Alternatively, specify a value such as other=-1.
	50	+ + data = tl.load(input + idx, mask = mask, care_padding=False) # Alternatively, specify a value such as other=-1.
	51	+ ```
	52	+
	53	+- Example 2: Using `for` loops in Triton operators to increase tiling and enhance DOP
	54	+
	55	+ In Triton operator programming, `mask` operations are frequently employed in syntax such as `load`, `store`, and `where`. During performance optimization, you should prioritize identifying performance degradation caused by these operations. When the logic within Triton operators executes sequentially in a single pass (Start -> Data-in -> Computation -> Data-out -> End), instructions cannot be parallelized, resulting in low execution efficiency. By introducing `for` loops to increase tiling, you can process data in multiple passes (with each pass handling a reduced data volume), enabling parallel execution of data-in, computation, and data-out. This approach reduces serial waiting time and improves overall performance. Additionally, compared to monolithic (non-tiled) data processing, the use of `for` loops for tiling reduces UB space consumption.
	56	+ Note: Mathematical equivalence is an important aspect to consider when you increase data tiling.
	57	+
	58	+ ```diff
	59	+ @triton.jit
	60	+ def alloc_extend_kernel(
	61	+ pre_lens_ptr,
	62	+ seq_lens_ptr,
	63	+ free_page_ptr,
	64	+ out_indices,
	65	+ bs_upper: tl.constexpr,
	66	+ page_size: tl.constexpr,
	67	+ max_num_extend_tokens: tl.constexpr,
	68	+ + BLOCK_SIZE: tl.constexpr = 1024,
	69	+ ):
	70	+ pid = tl.program_id(0)
	71	+
	72	+ load_offset = tl.arange(0, bs_upper)
	73	+ seq_lens = tl.load(seq_lens_ptr + load_offset, mask=load_offset <= pid)
	74	+ pre_lens = tl.load(pre_lens_ptr + load_offset, mask=load_offset <= pid)
	75	+ extend_lens = seq_lens - pre_lens
	76	+
	77	+ seq_len = tl.load(seq_lens_ptr + pid)
	78	+ pre_len = tl.load(pre_lens_ptr + pid)
	79	+ extend_len = seq_len - pre_len
	80	+
	81	+ sum_extend_lens = tl.sum(extend_lens)
	82	+ output_start_loc = sum_extend_lens - extend_len
	83	+
	84	+ num_pages_after = (seq_lens + page_size - 1) // page_size
	85	+ num_pages_before = (pre_lens + page_size - 1) // page_size
	86	+ num_new_pages = num_pages_after - num_pages_before
	87	+
	88	+ num_page_start_loc_self = (seq_len + page_size - 1) // page_size - (
	89	+ pre_len + page_size - 1
	90	+ ) // page_size
	91	+ sum_num_new_pages = tl.sum(num_new_pages)
	92	+ new_page_start_loc = sum_num_new_pages - num_page_start_loc_self
	93	+
	94	+ # Part 2: fill the new full pages
	95	+ num_part2 = (
	96	+ seq_len // page_size * page_size
	97	+ - (pre_len + page_size - 1) // page_size * page_size
	98	+ )
	99	+
	100	+ - # load data at once
	101	+ - offset_many_page = tl.arange(0, max_num_extend_tokens)
	102	+ - page_start = tl.load(
	103	+ - free_page_ptr + new_page_start_loc + offset_many_page // page_size,
	104	+ - mask=offset_many_page < num_part2,
	105	+ - )
	106	+ - tl.store(
	107	+ - out_indices + output_start_loc + offset_many_page,
	108	+ - page_start * page_size + offset_many_page % page_size,
	109	+ - mask=offset_many_page < num_part2,
	110	+ - )
	111	+
	112	+ + # load data using loop
	113	+ + num_loop = tl.cdiv(max_num_extend_tokens, BLOCK_SIZE)
	114	+ + blk_offset = tl.arange(0, BLOCK_SIZE)
	115	+ + for i in range(num_loop):
	116	+ + offset_many_page = blk_offset + i * BLOCK_SIZE
	117	+ + page_start = tl.load(
	118	+ + free_page_ptr + new_page_start_loc + offset_many_page // page_size,
	119	+ + mask=offset_many_page < num_part2,
	120	+ + )
	121	+ + tl.store(
	122	+ + out_indices + output_start_loc + offset_many_page,
	123	+ + page_start * page_size + offset_many_page % page_size,
	124	+ + mask=offset_many_page < num_part2,
	125	+ + )
	126	+ ```
	127	+
	128	+
	129	+## Optimizing Data Types
	130	+
	131	+### I. Core Principles of Data Type Optimization
	132	+
	133	+Some operations of the A2/A3 vector units do not support certain data types. In this case, the corresponding vector operations will degrade to scalar operations, affecting performance. If the overall operator accuracy is not affected, it is advisable to use supported data types to improve performance.
	134	+The following operations are involved.
	135	+\| Operator Name \| Unsupported Data Type \|
	136	+\|---\|---\|
	137	+\| Vector Add\| int64 \|
	138	+\| Vector Cmp\| int64/int32 \|
	139	+
	140	+### II. Code Examples
	141	+
	142	+- Example code of the Triton operator Vector Add
	143	+
	144	+ For the following Triton operator, when the input tensors `x` and `y` utilize the int64 data type, `x1 + y1` is expanded into a scalar operation, which degrades performance. Provided that computational accuracy remains unaffected, it is advisable to use the int32 data type.
	145	+ ``` diff
	146	+ @triton.jit
	147	+ def npu_vector_add_kernel(
	148	+ x, # [Tensor] input tensor (1 x col)
	149	+ y, # [Tensor] input tensor (1 x col)
	150	+ z, # [Tensor] output tensor (1 x col)
	151	+ vector_len: tl.constexpr, # len of the vector
	152	+ BLOCK_SIZE: tl.constexpr
	153	+ ):
	154	+ pid = tl.program_id(axis=0)
	155	+ offset = pid * BLOCK_SIZE + tl.arange(BLOCK_SIZE)
	156	+ len_mask = offset < vector_len
	157	+ x1 = tl.load(x + offset, mask=len_mask)
	158	+ y1 = tl.load(y + offset, mask=len_mask)
	159	+ z1 = x1 + y1
	160	+ tl.store(z + offset, z1, mask=len_mask)
	161	+ ```
	162	+
	163	+- Example code of the Triton operator Vector Cmp
	164	+
	165	+ In the following Triton operator, the `mask` operation utilizes Cmp. However, Cmp does not support the int64 or int32 data type, causing the condition `cols < N` to be expanded into a scalar operation, which reduces performance. Provided that computational accuracy remains unaffected, it is advisable to use the FP32 data type.
	166	+ In Triton operator programming, `mask` operations are frequently employed in syntax such as `load`, `store`, and `where`. During performance optimization, you should prioritize identifying performance degradation caused by these operations.
	167	+
	168	+ ``` diff
	169	+ @triton.jit
	170	+ def npu_vector_cmp_kernel(
	171	+ X, # [Tensor] input tensor (row x col)
	172	+ Out, # [Tensor] output tensor (row x col)
	173	+ Mean, # [Vector] mean tensor (row, ) of X
	174	+ Rstd, # [Vector] std tensor (row, ) of X
	175	+ stride_x_row, # [Scalar] stride of row of x
	176	+ stride_out_row, # [Scalar] stride of row of out, normally equals to stride_x_row
	177	+ M, # [Scalar] row number
	178	+ N, # [Scalar] col number
	179	+ eps, # [Scalar] epsilon to avoid division by zeros
	180	+ BLOCK_M: tl.constexpr,
	181	+ BLOCK_N: tl.constexpr
	182	+ ):
	183	+ group_m = tl.program_id(0)
	184	+ group_n = tl.program_id(1)
	185	+ row = group_m
	186	+
	187	+ # calculate index & offset
	188	+ Mean = Mean + group_n * M
	189	+ Rstd = Rstd + group_n * M
	190	+ X = X + row * stride_x_row + group_n * N
	191	+ Out = Out + row * stride_out_row + group_n * N
	192	+
	193	+ cols = tl.arange(0, BLOCK_N) # cols is int64
	194	+ x = tl.load(X + cols, mask=cols < N, other=0.0).to(tl.float32)
	195	+
	196	+ # calculate mean & rstd
	197	+ mean = tl.sum(x, axis=0) / N
	198	+ tl.store(Mean + row, mean)
	199	+ # [Changed begin]
	200	+ - xbar = tl.where(cols < N, X - mean, 0.0)
	201	+ + cols_cmp = cols.to(tl.float32)
	202	+ + xbar = tl.where(cols_cmp < N, x - mean, 0.0)
	203	+ # [Changed end]
	204	+
	205	+ var = tl.sum(xbar * xbar, axis=0) / N
	206	+ rstd = 1 / tl.sqrt(var + eps)
	207	+ tl.store(Rstd + row, rstd)
	208	+
	209	+ # calculate Out
	210	+ mask = cols < N
	211	+ out = (x - mean) * rstd
	212	+ tl.store(Out + cols, out, mask=mask)
	213	+ ```

Adocs/en/programming_guide.md+429-0

@@ -0,0 +1,429 @@
	1	+# Triton Operator Development Guide
	2	+This document focuses on the issues that need to be paid attention to during Triton operator development on NPUs, which are divided into three aspects: multi-core task parallelism, single-core data transfer, and single-core data computation. First, section "Multi-Core Task Parallelism" describes the basis for setting the maximum number of hardware cores and the specific implementation. Then, section "Single-Core Data Transfer" describes how to set the proper data block size in a loop, introduces the common optimization methods, and describes how to handle the UB overflow problem that may occur. Finally, section "Single-Core Data Computation" focuses on how to develop Triton operators and highlights the key points.
	3	+
	4	+## Multi-Core Task Parallelism
	5	+
	6	+### Setting the Maximum Number of Hardware Cores
	7	+
	8	+In a Triton operator, the grid is usually used for core allocation. For GPUs, it contains dozens or hundreds of core SMs. However, for the Ascend NPU platform, it contains dozens of AI Cores for computation.\
	9	+Although the runtime interface allows a maximum of 65,535 concurrent tasks to be delivered, the tasks that exceed the number of physical cores are delivered in a new round. If the Triton operator on the GPU is directly executed on the Ascend platform, a large number of tasks will introduce considerable overhead during core startup and initialization, affecting the operator performance.\
	10	+Therefore, the core allocation logic needs to be modified based on the Ascend platform features. The most recommended method is to fix the number of cores to the number of physical cores of the hardware and perform more detailed data block division within the cores.
	11	+
	12	+* For pure vector operators, the number of cores is equal to the number of vector cores.
	13	+* For CV fusion operators, the number of cores is equal to the number of cube cores (usually half of the number of vector cores). During operator execution, vector cores are called at a ratio of 1:2.
	14	+
	15	+You can obtain the number of vector cores and number of cube cores through the following interfaces:
	16	+
	17	+```python
	18	+import torch_npu
	19	+import triton.runtime.driver as driver
	20	+import torch_npu
	21	+
	22	+device = torch_npu.npu.current_device()
	23	+properties = driver.active.utils.get_device_properties(device)
	24	+vectorcore_num = properties["num_vectorcore"]
	25	+aicore_num = properties["num_aicore"]
	26	+
	27	+```
	28	+
	29	+According to the sample code, fix the number of cores, and then process task blocks in batches through an internal loop.
	30	+
	31	+```python
	32	+grid = (NUM_CORE ,)
	33	+_attn_fwd[grid](Q, K, V, M, Out, acc, scale......)
	34	+
	35	+@triton.jit
	36	+def _attn_fwd(Q, K, V, M, Out, acc, scale,
	37	+ ......
	38	+ Z: tl.constexpr, H: tl.constexpr,
	39	+ N_CTX: tl.constexpr,
	40	+ HEAD_DIM: tl.constexpr,
	41	+ BLOCK_M: tl.constexpr,
	42	+ BLOCK_N: tl.constexpr,
	43	+ STAGE: tl.constexpr
	44	+ ):
	45	+ # Calculate the total number of tasks and flatten the three-dimensional tasks (Z, H, M) into a one-dimensional total number of tasks.
	46	+ NUM_BLOCKS_M = N_CTX // BLOCK_M
	47	+ NUM_BLOCKS = NUM_BLOCKS_M * Z * H
	48	+
	49	+ # Each core selects the task to be processed based on its own identifier.
	50	+ pid = tl.program_id(0) # Unique ID of the current core.
	51	+ NUM_CORE = tl.num_programs(0) # Obtain the total number of cores that are started.
	52	+ # Loop rule: range(pid, NUM_BLOCKS, NUM_CORE) implements step-based task allocation.
	53	+ # - Start value (pid): Each core obtains tasks from its own ID to avoid task overlapping.
	54	+ # - Step length (NUM_CORE): The step is based on the total number of cores to ensure that tasks are evenly allocated to each core.
	55	+ for block_idx in range(pid, NUM_BLOCKS, NUM_CORE):
	56	+ # Calculate the data offset of each task.
	57	+ # [Core: Reverse restoration of one-dimensional task index to original multi-dimensional index.]
	58	+ # block_idx is the one-dimensional task index after flattening. The original dimensions are restored through integer division and remainder.
	59	+ # 1. Split the Z+H combined axis and M block axis.
	60	+ # - Exact division by NUM_BLOCKS_M: Extract the index (task_hz_idx) of the Z+H combined axis.
	61	+ # - Remainder of NUM_BLOCKS_M: Extract the block index (task_m_idx) of the M dimension.
	62	+ task_hz_idx = block_idx // NUM_BLOCKS_M
	63	+ task_m_idx = block_idx % NUM_BLOCKS_M
	64	+ # 2. Split the Z+H combined axis into the original Z axis and H axis.
	65	+ # - Exact division by H: Restore the Z axis index (off_z).
	66	+ # - Remainder of H: Restore the H axis index (off_h).
	67	+ off_z = task_hz_idx // H
	68	+ off_h = task_hz_idx % H
	69	+ # 3. Calculate the data offset: Locate the start position of the corresponding data in the Q/K/V tensor based on the restored Z/H index.
	70	+ qvk_offset = off_z.to(tl.int64) * stride_qz + off_h.to(tl.int64) * stride_qh
	71	+```
	72	+
	73	+
	74	+## Single-Core Data Transfer
	75	+
	76	+### Setting the Proper Data Block Size (BLOCK SIZE)
	77	+Take add_kernel as an example. The variables and operations determine the on-chip memory usage. You can change the value of BLOCK_SIZE to adjust the size of the data block in the loop and the size of the intermediate result. If the upper limit is exceeded, the expected usage size is displayed and an error is reported during operator compilation. To achieve the maximum compute-to-memory ratio, BLOCK_SIZE needs to be as large as possible without exceeding the on-chip space. You can set different BLOCK_SIZE values in advance by using [autotune](#triton-autotune) of Triton-Ascend. The optimal setting is automatically selected during running.
	78	+
	79	+```python
	80	+@triton.jit
	81	+def add_kernel(x_ptr,
	82	+ y_ptr,
	83	+ out_ptr,
	84	+ n, # Total number of elements.
	85	+ BLOCK_SIZE: tl.constexpr, # Number of block elements.
	86	+ ):
	87	+ pid = tl.program_id(0)
	88	+ NUM_CORE = tl.num_programs(0)
	89	+ NUM_BLOCKS = tl.cdiv(n, BLOCK_SIZE)
	90	+ for block_idx in range(pid, NUM_BLOCKS, NUM_CORE):
	91	+ block_start = block_idx * BLOCK_SIZE
	92	+ # The block size is BLOCK_SIZE.
	93	+ offsets = block_start + tl.arange(0, BLOCK_SIZE)
	94	+ mask = offsets < n
	95	+ # Load data of x and y to the on-chip memory.
	96	+ x = tl.load(x_ptr + offsets, mask=mask)
	97	+ y = tl.load(y_ptr + offsets, mask=mask)
	98	+
	99	+ output = x + y
	100	+
	101	+ tl.store(output_ptr + offsets, output, mask=mask)
	102	+```
	103	+
	104	+### Aligning the Size of the Tail Axis of the Tensor
	105	+
	106	+[Description] For VV operators, if the Vector core needs to be called for computation, the UB of the Ascend hardware requires that the size of the tail axis of the tensor be divisible by 32 bytes. For CV operators, if the Vector core and Cube core need to be called for computation, the size of the tail axis of the tensor must be divisible by 512 bytes. If the tail axis length is insufficient, the tail axis length will be automatically padded. Under this premise, the performance of operations with the shape of (2048,3) and (2048,1) tensors in the model deteriorates significantly due to automatic padding. In this case, you can perform the transpose operation to convert the alignment axis to a lower dimension until the store operation is performed, avoiding automatic padding and optimizing the computing speed. In addition, the transpose operation is also affected by the automatic padding rule. Therefore, special skills are required to avoid padding. The following is a tip for "borrowing axis for transpose", which is applicable to the scenario where tensor.numel() % 256Byte == 0:
	107	+
	108	+- Note: VV operators indicate that only Vector Core is used during operator computation. CV operators indicate that both AI Core and Vector Core are used during operator computation.
	109	+- Example
	110	+
	111	+```python
	112	+# conv_state = tensor([2048, 3], bfloat16)
	113	+conv_state = tl.load(conv_state_ptr + conv_batch_offs * conv_batch_stride + doffs * 3 + tl.arange(0, 2048 * 3)) # It is considered as the 1D tensor load. In this case, numel is aligned and no padding is performed.
	114	+conv_state_T = conv_state.reshape(128, 16 * 3).trans().reshape(16, 3 * 128).trans().reshape(3 * 2048,) # The long axis (2048) is split into an aligned axis (16) and lent to the short axis (3) to align the two axes.
	115	+```
	116	+
	117	+### Transferring Data to the UB and Then Selecting the Target Value from the UB
	118	+
	119	+[Description] In the discrete scenario of the NPU, data can be transferred to the UB and then the target value can be selected from share.
	120	+
	121	+- Example
	122	+
	123	+```diff
	124	+@triton.jit
	125	+def pick_kernel(
	126	+ x_ptr,
	127	+ idx_ptr,
	128	+ y_ptr,
	129	+ stride_x,
	130	+ stride_idx,
	131	+ stride_y,
	132	+ M: tl.constexpr,
	133	+ N: tl.constexpr
	134	+):
	135	+ pid = tl.program_id(0)
	136	+ rn = tl.arange(0, N)
	137	+
	138	+ idx = tl.load(idx_ptr + rn * stride_idx)
	139	+ mask = idx < M
	140	+
	141	+- val = tl.load(x_ptr + idx * stride_x, mask=mask)
	142	+
	143	++ rm = tl.arange(0, M)
	144	++ x_shared = tl.load(x_ptr + rm * stride_x) # [M]
	145	++ val = tl.gather(x_shared, idx, 0)
	146	+
	147	+ tl.store(y_ptr + rn * stride_y, val, mask=mask)
	148	+```
	149	+
	150	+- Performance analysis and comparison before and after optimization
	151	+
	152	+You can use the msProf tool to execute the test case to obtain the PROF_\* folder, which contains the op_summary_\*.csv file. This file can be used to analyze the pipeline. Note: \* indicates the timestamp. For details, see the [performance data collection methods](./debug_guide/profiling.md).
	153	+
	154	+\|\|Op Name\|aiv_time(us)\|
	155	+\|:---- \|:--------\|:--------\|
	156	+\|Unoptimized\|pick_kernel\|574.124\|
	157	+\|Optimized\|pick_kernel\|171.175\|
	158	+
	159	+According to the data in the table, the values of aiv_mte2_time(us) and aiv_mte2_ratio before and after the optimization are greatly different. The optimization solution first transfers most of the data to the UB, reducing the number of times that small batches of data are transferred to the UB through the L2 and the total time for transferring data to the UB through the L2.
	160	+
	161	+### Parallel Storage and Computation
	162	+
	163	+Triton-Ascend supports two data processing modes: serial storage and computation and parallel storage and computation.
	164	+
	165	+Serial storage and computation: Data is first transferred from the global memory to the on-chip memory, and then the next batch of data is transferred after the computation is complete. This mode has a significant idle waiting time, and the efficiency is low.
	166	+
	167	+Parallel storage and computation: Computing is performed when the first batch of data is transferred to the on-chip memory. Then, the second batch of data is transferred, and the "transfer + compute" pipeline operation is formed, significantly improving the overall throughput.
	168	+
	169	+The key to implementing parallel storage and computation is to properly design the data tiling policy so that the data required for the next phase can be prepared in advance during the compute of the current batch of data, thereby implementing parallelization of data transfer and computing.
	170	+ Currently, the compiler is configured with multiBuffer set to True by default, and the parallel storage and computation are supported by default.
	171	+
	172	+### Tiling Optimization
	173	+
	174	+Before the AI Core performs computation, data needs to be transferred to the on-chip memory. The on-chip memory space is usually much smaller than the total data volume to be processed by the AI Core. For example, the on-chip memory capacity of Atlas 800T/I A2 is 192 KB. After doublebuffer is enabled by default, the capacity is reduced to half of the original capacity. Therefore, data needs to be tiled during operator computation, and only a small part of the data is loaded and processed each time.
	175	+
	176	+- Example
	177	+
	178	+```diff
	179	+@libentry()
	180	+@triton.autotune(configs=runtime.get_tuned_config("masked_fill"), key=["N"])
	181	+@triton.jit
	182	+- def masked_fill_kernel(inp, expand_mask, value, out, N, BLOCK_SIZE: tl.constexpr):
	183	++ def masked_fill_kernel(inp, expand_mask, value, out, N, BLOCK_SIZE: tl.constexpr, BLOCK_SIZE_SUB: tl.constexpr):
	184	+ pid = tl.program_id(axis=0)
	185	++ base_offset = pid * BLOCK_SIZE
	186	+
	187	++ # Calculate the total number of blocks that need to be processed
	188	++ num_sub_blocks = BLOCK_SIZE // BLOCK_SIZE_SUB
	189	+
	190	++ # Loop processing each sub block
	191	++ for sub_block_idx in range(num_sub_blocks):
	192	++ # Calculate the offset of the current sub block
	193	++ sub_offset = base_offset + sub_block_idx * BLOCK_SIZE_SUB
	194	++ offsets = sub_offset + tl.arange(0, BLOCK_SIZE_SUB)
	195	+- offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
	196	+ mask = offsets < N
	197	+ # Load input and mask
	198	+ input_vals = tl.load(inp + offsets, mask=mask, other=0)
	199	+ fill_mask_vals = tl.load(expand_mask + offsets, mask=mask, other=0).to(tl.int1)
	200	+
	201	+ # Write the original input first
	202	+ tl.store(out + offsets, input_vals, mask=mask)
	203	+
	204	+ # Overlay and write value at the position that needs to be filled
	205	+- value_to_write = tl.full([BLOCK_SIZE], value, dtype=input_vals.dtype)
	206	++ value_to_write = tl.full([BLOCK_SIZE_SUB], value, dtype=input_vals.dtype)
	207	+ overwrite_vals = tl.where(fill_mask_vals, value_to_write, tl.load(out + offsets, mask=mask, other=0))
	208	+ tl.store(out + offsets, overwrite_vals, mask=mask)
	209	+```
	210	+
	211	+### Triton Autotune
	212	+
	213	+In tiling block optimization, the values of block parameters such as BLOCK_SIZE and BLOCK_SIZE_SUB directly affect the operator performance. However, manually debugging parameter combinations is inefficient and it is difficult to find the optimal values. triton.autotune is an autotune tool provided by the Triton framework. It can traverse preset parameter configurations, compare the performance of different parameter configurations, and automatically select the optimal parameter combination. It is the core auxiliary means of tiling optimization.
	214	+
	215	+- Core functions
	216	+Automatically traversal of the parameter space: Test the performance of different values of block parameters of the constexpr type such as BLOCK_SIZE and BLOCK_SIZE_SUB in batches.
	217	+Performance benchmark comparison: Select the optimal parameters that adapt to the current hardware based on the operator execution duration.
	218	+Tuning result caching: The optimal configuration after tuning is cached. The optimal configuration is reused when the operator is called, avoiding repeated tuning.
	219	+
	220	+- Simple example
	221	+ ```diff
	222	+ @triton.autotune(
	223	+ configs=[ # List of parameter configurations to be tested. The candidate parameter values must be powers of 2.
	224	+ triton.Config({'BLOCK_SIZE': 128}),
	225	+ triton.Config({'BLOCK_SIZE': 256}),
	226	+ triton.Config({'BLOCK_SIZE': 512}),
	227	+ ],
	228	+ key=['n_elements'], # Tune dimension: input dimension on which the parameter value depends.
	229	+ )
	230	+ @triton.jit
	231	+ def add_kernel(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE: tl.constexpr):
	232	+ pid = tl.program_id(axis=0)
	233	+ block_start = pid * BLOCK_SIZE
	234	+ offsets = block_start + tl.arange(0, BLOCK_SIZE)
	235	+ mask = offsets < n_elements
	236	+
	237	+ x = tl.load(x_ptr + offsets, mask=mask)
	238	+ y = tl.load(y_ptr + offsets, mask=mask)
	239	+ output = x + y
	240	+ tl.store(output_ptr + offsets, output, mask=mask)
	241	+ ```
	242	+- Note: You can set the following environment variables to print the optimal parameter information.
	243	+ ```diff
	244	+ export TRITON_PRINT_AUTOTUNING=1
	245	+ ```
	246	+
	247	+### How Do I Avoid UB Overflow on the NPU?
	248	+
	249	+[Description] On the NPU, the UB or L1 size has an upper limit. When this error occurs, reduce the amount of data transferred at a time and use the for loop to process long sequences.
	250	+```diff
	251	+E triton.compiler. errors.MLIRCompilationError:
	252	+E ///--------------------- [ERROR][Triton][BEG]-------------------------
	253	+E [ConvertLinalgRToBinary] encounters error:
	254	+E loc("/tmp/tmpsb6qkdih/kernel.ttadapter.mlir":2:1): error: Failed to run BishengHIR pipeline
	255	+E
	256	+E loc("/tmp/tmpsb6qkdih/kernel.ttadapter.mlir":3:3): error: ub overflow, requires 3072256 bits while 1572864 bits available! (possible reason
	257	+large or block number is more than what user expect due to multi-buffer feature is enabled and some ops need extra local buffer. )
	258	+```
	259	+[Note] The UB size of the A2 series products is 192 KB (1,572,864 bits).
	260	+
	261	+
	262	+## Single-Core Data Computation
	263	+### R&D Goals
	264	+Implement basic data operation operators (such as addition, subtraction, multiplication, division, activation functions, and simple matrix element operations) on the Ascend NPU single core. Ensure that operators are efficiently executed on a single core, laying a foundation for subsequent multi-core parallel processing and distributed expansion.
	265	+
	266	+
	267	+### Development Procedure
	268	+1. Determine the operator function.
	269	+-Determine the shapes and data types (such as float16, float32, and int32) of the input and output tensors.
	270	+-Check whether broadcast and boundary processing are required.
	271	+
	272	+
	273	+2. Write kernel functions.
	274	+Single-kernel computation corresponds to block-level data processing.
	275	+Single-kernel data computation example: vector addition
	276	+```diff
	277	+import torch
	278	+import torch_npu
	279	+
	280	+import triton
	281	+import triton.language as tl
	282	+
	283	+@triton.jit
	284	+def add_kernel(x_ptr, # Pointer to first input vector.
	285	+ y_ptr, # Pointer to second input vector.
	286	+ output_ptr, # Pointer to output vector.
	287	+ n_elements, # Size of the vector.
	288	+ BLOCK_SIZE: tl.constexpr, # Number of elements each program should process.
	289	+ # NOTE: constexpr so it can be used as a shape value.
	290	+):
	291	+ pid = tl.program_id(axis=0) # We use a 1D launch grid so axis is 0.
	292	+ block_start = pid * BLOCK_SIZE
	293	+ offsets = block_start + tl.arange(0, BLOCK_SIZE)
	294	+ mask = offsets < n_elements
	295	+ x = tl.load(x_ptr + offsets, mask=mask)
	296	+ y = tl.load(y_ptr + offsets, mask=mask)
	297	+ output = x + y
	298	+ tl.store(output_ptr + offsets, output, mask=mask)
	299	+```
	300	+Calling:
	301	+ ```diff
	302	+def add(x: torch.Tensor, y: torch.Tensor):
	303	+ output = torch.empty_like(x)
	304	+ n_elements = output.numel()
	305	+ grid = lambda meta: (triton.cdiv(n_elements, meta['BLOCK_SIZE']), )
	306	+ add_kernel[grid](x, y, output, n_elements, BLOCK_SIZE=1024)
	307	+ return output
	308	+```
	309	+Use the above function to compute element-wise sum of two torch.tensor objects and test its correctness.
	310	+ ```diff
	311	+torch.manual_seed(0)
	312	+size = 98432
	313	+x = torch.rand(size, device='npu')
	314	+y = torch.rand(size, device='npu')
	315	+output_torch = x + y
	316	+output_triton = add(x, y)
	317	+print(output_torch)
	318	+print(output_triton)
	319	+print(f'The maximum difference between torch and triton is '
	320	+f'{torch.max(torch.abs(output_torch - output_triton))}')
	321	+# Out:
	322	+# tensor([1.3713, 1.3076, 0.4940, ..., 0.6724, 1.2141, 0.9733], device='npu')
	323	+# tensor([1.3713, 1.3076, 0.4940, ..., 0.6724, 1.2141, 0.9733], device='npu')
	324	+# The maximum difference between torch and triton is 0.0
	325	+```
	326	+
	327	+
	328	+3. Key points of single-kernel computation
	329	+
	330	+-Block-level data processing: Each computing block is responsible for a small segment of data, ensuring parallelism.
	331	+
	332	+-Boundary check: Use mask or if (tid < N) to avoid out-of-bounds access.
	333	+
	334	+-Block size selection: Properly set the block and grid.
	335	+
	336	+
	337	+4. Performance points
	338	+(1) Memory access optimization
	339	+-Ensure sequential access.
	340	+-Use the aligned stride to avoid cross-row/cross-column jump access.
	341	+-Align the data block size to the 32-byte boundary.
	342	+Ensure that the input and output buffers are aligned during allocation to avoid memory access performance deterioration.
	343	+Example:
	344	+ ```diff
	345	+BLOCK_SIZE = 256 # 256 x 4 bytes = 1024 bytes, which are well-aligned.
	346	+
	347	+@triton.jit
	348	+def vec_add_kernel(X, Y, Z, N,
	349	+ BLOCK_SIZE: tl.constexpr):
	350	+ pid = tl.program_id(axis=0)
	351	+
	352	+ # Compute the index range of the current block.
	353	+ offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
	354	+
	355	+ # The mask is used to prevent out-of-bounds.
	356	+ mask = offsets < N
	357	+
	358	+ # Contiguous memory access: The offsets are contiguous.
	359	+ x = tl.load(X + offsets, mask=mask)
	360	+ y = tl.load(Y + offsets, mask=mask)
	361	+
	362	+ z = x + y
	363	+
	364	+ # Contiguous writeback
	365	+ tl.store(Z + offsets, z, mask=mask)
	366	+
	367	+
	368	+def vec_add(x, y):
	369	+ assert x.numel() == y.numel()
	370	+ N = x.numel()
	371	+
	372	+ # Allocate aligned memory. (PyTorch is aligned to 64 bytes by default.)
	373	+ z = torch.empty_like(x)
	374	+
	375	+ # grid: Each block processes BLOCK_SIZE elements.
	376	+ grid = lambda meta: (triton.cdiv(N, meta['BLOCK_SIZE']),)
	377	+
	378	+ vec_add_kernel[grid](x, y, z, N, BLOCK_SIZE=BLOCK_SIZE)
	379	+
	380	+ return z
	381	+```
	382	+
	383	+(2) Sub-block division
	384	+-Divide a large matrix into small blocks. Each block is computed in the UB.
	385	+-Sub-block division should ensure both memory access continuity and computing unit utilization.
	386	+Example:
	387	+ ```diff
	388	+BLOCK_M = 64 # Each block processes 64 rows.
	389	+BLOCK_N = 64 # Each block processes 64 columns.
	390	+BLOCK_K = 32 # Internal dimension is accumulated.
	391	+
	392	+@triton.jit
	393	+def matmul_kernel(
	394	+ A, B, C,
	395	+ M, N, K,
	396	+ stride_am, stride_ak,
	397	+ stride_bk, stride_bn,
	398	+ stride_cm, stride_cn,
	399	+ BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr
	400	+):
	401	+ pid_m = tl.program_id(0) # ID of the block in the M direction.
	402	+ pid_n = tl.program_id(1) # ID of the block in the N direction.
	403	+
	404	+ # Start coordinates of the current block.
	405	+ offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
	406	+ offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
	407	+ offs_k = tl.arange(0, BLOCK_K)
	408	+
	409	+ # Initialize accumulators.
	410	+ acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
	411	+
	412	+ # Compute blocks in the loop.
	413	+ for k in range(0, K, BLOCK_K):
	414	+ a = tl.load(
	415	+ A + (offs_m[:, None] * stride_am + (offs_k[None, :] + k) * stride_ak),
	416	+ mask=(offs_m[:, None] < M) & (offs_k[None, :] + k < K),
	417	+ other=0.0
	418	+ )
	419	+ b = tl.load(
	420	+ B + ((offs_k[:, None] + k) * stride_bk + offs_n[None, :] * stride_bn),
	421	+ mask=(offs_k[:, None] + k < K) & (offs_n[None, :] < N),
	422	+ other=0.0
	423	+ )
	424	+ acc += tl.dot(a, b)
	425	+
	426	+ # Write back the result.
	427	+ c = C + (offs_m[:, None] * stride_cm + offs_n[None, :] * stride_cn)
	428	+ tl.store(c, acc, mask=(offs_m[:, None] < M) & (offs_n[None, :] < N))
	429	+```

Adocs/en/quick_start.md+60-0

@@ -0,0 +1,60 @@
	1	+# Quick Start
	2	+
	3	+## Project Overview
	4	+
	5	+Triton-Ascend is an optimized version of Triton that adapts to Huawei Ascend chips. It provides efficient automatic optimization of kernel functions, operator compilation, and deployment capabilities, and supports products such as Ascend Atlas A2/A3.
	6	+While being compatible with the core syntax of Triton, Ascend is optimized for features of Ascend NPUs, including automatic parsing of kernel function parameters, memory access logic optimization, and security deployment mechanism optimization.
	7	+
	8	+## Online Documents
	9	+Complete online documents and network materials are provided, covering environment setup, operator development, optimization practices, and FAQ, to help you get started quickly. For details, see the [online documents](https://triton-ascend.readthedocs.io/zh-cn/latest/index.html).
	10	+
	11	+## Environment Requirements
	12	+### Hardware Requirements
	13	+Supported OS: Linux (AArch64/x86_64)
	14	+
	15	+Supported Ascend products: Atlas A2/A3 series
	16	+
	17	+Minimum hardware configuration: single-device 32 GB graphics memory (recommended)
	18	+
	19	+### Software Dependency
	20	+Python (Python 3.9 to Python 3.11), CANN_TOOLKIT, CANN_OPS, [requirements.txt](../../requirements.txt), and [requirements_dev.txt](../../requirements_dev.txt)
	21	+
	22	+For details about the CANN installation and configuration script, see [CANN installation description](https://www.hiascend.com/document/detail/zh/canncommercial/850/softwareinst/instg/instg_0000.html?Mode=PmIns&InstallType=local&OS=Ubuntu). The quick installation commands are as follows:
	23	+```bash
	24	+chmod +x Ascend-cann-toolkit_8.5.0_linux-aarch64.run
	25	+chmod +x Ascend-cann-A3-ops_8.5.0_linux-aarch64.run
	26	+
	27	+sudo ./Ascend-cann-toolkit_8.5.0_linux-aarch64.run --install
	28	+sudo ./Ascend-cann-A3-ops_8.5.0_linux-aarch64.run --install
	29	+```
	30	+
	31	+- Note: [CANN_TOOLKIT and CANN_OPS](https://www.hiascend.com/developer/download/community/result?module=cann&cann=8.5.0) are the key tool packages for enabling the Ascend computing card.
	32	+You need to select the required version (8.5.0 is recommended) based on the Ascend card model you use. The CANN installation takes about 5 to 10 minutes. Wait until the installation is complete.
	33	+
	34	+You can run the following command to install the requirements:
	35	+```shell
	36	+pip install -r requirements.txt -r requirements_dev.txt
	37	+```
	38	+
	39	+## Environment Setup
	40	+You can set up the Triton-Ascend environment by referring to section "Preparing the Environment" in [Installation Guide](installation_guide.md).
	41	+
	42	+### Obtaining the Triton-Ascend Software Package
	43	+You can install the latest stable version package using the CLI.
	44	+```shell
	45	+pip install triton-ascend
	46	+```
	47	+You can also download the nightly package from the [download link](https://test.pypi.org/project/triton-ascend/#history) and install it locally.
	48	+
	49	+- Note 1: If you download the nightly package for installation, select the Python version and architecture (AArch64/x86_64) of your server when selecting the Triton-Ascend package.
	50	+- Note 2: The nightly package is built every day. Developers submit MRs frequently. Note that if the package does not pass the stable test, function bugs may exist.
	51	+
	52	+## Example for Running Triton
	53	+
	54	+Run the [01-vector-add.py](../../ascend/examples/tutorials/01-vector-add.py) instance.
	55	+```bash
	56	+# Set the CANN environment variables (for example, as the root user and with the default installation path /usr/local/Ascend).
	57	+source /usr/local/Ascend/ascend-toolkit/set_env.sh
	58	+# Run the tutorials example.
	59	+python3 ./triton-ascend/ascend/examples/tutorials/01-vector-add.py
	60	+```

Adocs/en/release_note.md+50-0

@@ -0,0 +1,50 @@
	1	+# Triton-Ascend Release
	2	+
	3	+The Triton-Ascend version provides a stable code base snapshot, which is encapsulated into a binary package that can be easily installed through PyPI. In addition, the release represents that the development team can officially announce the availability of new functions, completed improvements, and changes that may affect users (such as destructive changes) to the community.
	4	+
	5	+## Release Compatibility Matrix
	6	+
	7	+The release compatibility matrix of the Triton-Ascend version is as follows.
	8	+
	9	+\| Triton-Ascend Version\| Python Version\| Manylinux Version\| Hardware Platform\| Hardware Product\|
	10	+\| --- \| --- \| --- \| --- \| --- \|
	11	+\| 3.2.0 \| 3.9 to 3.11\| glibc 2.27+, x86-64, AArch64 \| Ascend NPU \| Atlas A2/A3\|
	12	+
	13	+## Release Date
	14	+
	15	+The following is the release plan of Triton-Ascend. Note: The patch version is optional.
	16	+
	17	+\| Major Version\| Release Branch Cut-Out Time\| Release Date\| Patch Release Date\|
	18	+\| --- \| --- \| --- \| --- \|
	19	+\| 3.2.0 \| 2025-12-08\| 2026-01\| --- \|
	20	+
	21	+## Highlights
	22	+
	23	+### Triton-Ascend 3.2.0
	24	+
	25	+First release: Ascend NPU is supported.
	26	+
	27	+Triton-Ascend 3.2.0 is the first Triton version that officially supports Huawei Ascend NPU. This version is based on the Triton 3.2.0 community version and is specially adapted to the Ascend NPU hardware architecture.
	28	+
	29	+#### Main Features
	30	+
	31	+1. Full-stack support for Ascend NPU
	32	+ - The instruction set compilation pipeline from Triton IR to NPU is complete.
	33	+ - All Triton Ops are supported.
	34	+
	35	+2. Performance optimization
	36	+ - NPU-specific kernel optimization
	37	+ - CV compute optimization
	38	+
	39	+3. Developer tools
	40	+ - Comprehensive debug output is supported.
	41	+ - Intermediate compilation products are dumped.
	42	+
	43	+#### Known Limitations
	44	+
	45	+1. Data type: Some data types are still being improved.
	46	+2. Operator coverage: The supported operator set is being continuously expanded.
	47	+
	48	+#### Migration Guide
	49	+
	50	+For details about how to migrate existing Triton GPU users to Ascend NPU, see [Migrating Triton Operators from GPUs](./migration_guide/migrate_from_gpu.md).

Adocs/en/requirement.txt+3-0

@@ -0,0 +1,3 @@
	1	+sphinx
	2	+sphinx_rtd_theme
	3	+myst_parser

Rdocs/en/keep.md→docs/en/triton_api_extention/index.md+0-0

文件重命名但无更改。

@@ -0,0 +1,35 @@
	1	+## Environment Variables
	2	+
	3	+The following table describes how to set environment variables.
	4	+
	5	+\| Category\| Environment Variable\| Default Value\| Function Description\| Setting Description\| Change Description\|
	6	+\|------\|----------\|--------\|----------\|----------\|----------\|
	7	+\| Debugging and logging\| TRITON_DEBUG \| 0 or not set\| Specifies whether to enable the debugging output function of Triton to print detailed debugging information during running. This is useful for troubleshooting problems in the compilation or execution phase. When this parameter is set to 1, Triton outputs more information about the compilation, kernel generation, and execution. Some implementations may support more fine-grained debugging levels (such as 2 and 3), depending on the Triton version and implementation.\| 0: The debugging is disabled.<br>1: The debugging is enabled.\| \|
	8	+\| Debugging and logging\| MLIR_ENABLE_DUMP \| 0 or not set\| Specifies whether to dump the intermediate representation (IR) of all kernels before each MLIR optimization. You can set `MLIR_ENABLE_DUMP` to `kernelName` to dump the IR of a specific kernel.\| 0: Do not dump.<br>1: Dump the IR of all kernels.<br>kernelName: Dump the IR of a specific kernel.\| The Triton cache may interfere with the dump. If `MLIR_ENABLE_DUMP=1` does not take effect, you can run `rm -r ~/.triton/cache/*` to clear the Triton cache.\|
	9	+\| Debugging and logging\| LLVM_IR_ENABLE_DUMP \| 0 or not set\| Specifies whether to dump the IR before each LLVM IR optimization.\| 0: Do not dump.<br>1: Dump IRs.\| \|
	10	+\| Debugging and logging\| TRITON_REPRODUCER_PATH \| Not set\| Generates the MLIR reproduction file before each MLIR compilation phase. If a phase fails, `<reproducer_path>` saves the MLIR status before the failure.\| `<reproducer_path>`: save path.\| \|
	11	+\| Debugging and logging\| TRITON_INTERPRET \| 0 or not set\| Specifies whether to use the Triton interpreter instead of the GPU for running and support inserting Python breakpoints in kernel function code.\| 0: Breakpoints are not supported.<br>1: Breakpoints are supported.\| \|
	12	+\| Debugging and logging\| TRITON_ENABLE_LLVM_DEBUG \| 0 or not set\| Specifies whether to pass the`-debug` parameter to LLVM and outputs a large amount of debugging information. If there is too much information, you can use `TRITON_LLVM_DEBUG_ONLY` to limit the output scope.\| 0: Pass.<br>1: Do not pass.\| Another method to reduce output interference is as follows: Set the running program by setting `LLVM_IR_ENABLE_DUMP` to `1`, extract the IR before the target LLVM optimization channel, and run the `opt` tool of the LLVM separately. In this case, you can add `-debug-only=foo` to the command line to limit the debugging range.\|
	13	+\| Debugging and logging\| TRITON_LLVM_DEBUG_ONLY \| Not set\| Equivalent to the `-debug-only` command line option of LLVM. This parameter can be used to limit the LLVM debugging output to a specific optimization channel or component name (defined by the `#define DEBUG_TYPE` macro in LLVM and Triton), thereby effectively reducing redundant debugging output. You can specify one or more comma-separated values, for example, `TRITON_LLVM_DEBUG_ONLY="tritongpu-remove-layout-conversions"` or `TRITON_LLVM_DEBUG_ONLY="tritongpu-remove-layout-conversions,regalloc"`.\| Comma-separated values: channel or component name\| \|
	14	+\| Debugging and logging\| USE_IR_LOC \| 0 or not set\| Specifies whether to include location information (such as file names and line numbers) in the generated IR. This information is helpful for debugging, but may increase the size of the generated IR. If this parameter is set to 1, the IR is re-parsed, and the location information is mapped to the line number of the IR file with a specific extension (not the line number of the Python source file). This enables a direct mapping from the IR to the LLVM IR/PTX. When used with the performance analysis tool, this parameter can be used to implement fine-grained performance analysis on IR instructions.\| 0: No location information is included.<br>1: The location information is included.\| \|
	15	+\| Debugging and logging\| TRITON_PRINT_AUTOTUNING \| 0 or not set\| After the automatic optimization is complete, the optimal configuration and total time of each kernel are output.\| 0: Do not output.<br>1: Output.\| \|
	16	+\| Debugging and logging\| MLIR_ENABLE_REMARK \| 0 or not set\| Specifies whether to enable the output of remarks during MLIR compilation, including performance warnings in remarks.\| 0: Disabled.<br>1: Enabled.\| \|
	17	+\| Debugging and logging\| TRITON_KERNEL_DUMP \| 0 or not set\| Specifies whether to enable the dump function of the Triton kernel. When this function is enabled, Triton saves the generated kernel code (IR and final PTX in each compilation phase) to the specified directory.\| 0: Disabled.<br>1: Enabled.\| \|
	18	+\| Debugging and logging\| TRITON_DUMP_DIR \| Current working directory or not set\| Specifies the directory for storing the Triton kernel dump file, which is the directory for saving the IR and PTX when `TRITON_KERNEL_DUMP` is set to `1`.\| "path": save path.\| \|
	19	+\| Debugging and logging\| TRITON_DEVICE_PRINT \| 0 or not set\| If this parameter is set to `1` or `true` (`TRUE` is converted to `true`), the function of `tl.device_print` is enabled. Note: This function uses the GM buffer (the pointer of which is passed to the kernel).\| 0: Disabled.<br>1: The functionality of `tl.device_print` is enabled.\| The maximum size of the GM buffer for each thread is 16 KB. If the buffer size exceeds 16 KB, the excess content will be discarded. The value is fixed currently and will be adjusted through an environment variable.\|
	20	+\| Compilation control\| TRITON_ALWAYS_COMPILE \| 0 or not set\| Specifies whether Triton forcibly recompiles the kernel each time it runs, instead of using the existing cached version. By default, Triton caches the compiled kernels (based on parameters and configurations) to improve performance. If this parameter is set to 1, Triton ignores the cache and recompiles the kernel each time it runs, which is useful for debugging or testing new compiler features.\| 0: Disabled.<br>1: All kernels are recompiled during each running.\| \|
	21	+\| Compilation control\| DISABLE_LLVM_OPT \| 0 or not set\| If this parameter is set to 1, the optimization steps (LLVM optimization of make_llir and make_ptx) during LLVM compilation can be disabled. If this parameter is set to a character string, the LLVM optimization flags to be disabled are parsed. For example, if `DISABLE_LLVM_OPT` is set to `"disable-lsr"`, the loop strength optimization is disabled (this optimization may cause a performance fluctuation of up to 10% in some kernels with register pressure).\| 0: The LLVM optimization is enabled.<br>1: The optimization steps (LLVM optimization of make_llir and make_ptx) during LLVM compilation are disabled.\| \|
	22	+\| Compilation control\| MLIR_ENABLE_TIMING \| 0 or not set\| Specifies whether to enable the time statistics function during MLIR compilation.\| 0: Disabled.<br>1: Enabled.\| \|
	23	+\| Compilation control\| LLVM_ENABLE_TIMING \| 0 or not set\| Specifies whether to enable the time statistics function during LLVM compilation.\| 0: Disabled.<br>1: Enabled.\| \|
	24	+\| Compilation control\| TRITON_DEFAULT_FP_FUSION \| 1 (enabled)\| Specifies whether to enable the floating-point operation fusion optimization by default. The default floating-point operation fusion behavior (for example, mul+add->fma) is overwritten.\| 0: Disabled.<br>1: Enabled.\| \|
	25	+\| Compilation control\| TRITON_KERNEL_OVERRIDE \| 0 or not set\| Specifies whether to enable the Triton kernel override function. You can use the user-specified external file (such as IR/PTX) to override the default generated kernel code at the beginning of each compilation phase.\| 0: Disabled.<br>1: Enabled.\| \|
	26	+\| Compilation control\| TRITON_OVERRIDE_DIR \| Current working directory or not set\| Specifies the directory for searching the Triton kernel override file. Directory for loading the IR/PTX file when `TRITON_KERNEL_OVERRIDE` is set to `1`.\| "path": save path.\| \|
	27	+\| Compilation control\| TRITON_ASCEND_COMPILE_SPEED_OPT \| 0 or not set\| Specifies whether the JIT compiler skips the subsequent compilation phase after detecting that the kernel compilation fails. Set the parameter to `1` to skip the attempt. (The default value `0` indicates that the attempt is continued.)\| 0: Continue the attempt.<br>1: Skip.\| \|
	28	+\| Compilation control\| TRITON_COMPILE_ONLY \| 0 or not set\| Specifies whether to perform only compilation without execution. This parameter is used when remote_launch is used.\| 0: The optimization is disabled.<br>1: Enabled.\| \|
	29	+\| Compilation control\| TRITON_DISABLE_FFTS \| 0 or not set\| Specifies whether to enable FFTS.\| 0: The optimization is disabled.<br>1: Enabled.\| \|
	30	+\| Running and scheduling\| TRITON_ALL_BLOCKS_PARALLEL \| 0 or not set\| Specifies whether to enable the automatic optimization of the number of logical cores based on the number of physical cores. This parameter can be enabled only when logical cores can execute in parallel. When the number of logical cores is greater than the number of physical cores, enabling this parameter will instruct the compiler to automatically adjust the number of logical cores to match the number of physical cores, thereby reducing scheduling overhead. After this parameter is enabled, the value of grid can be greater than 65535. Limitation: This option can be enabled only when the logic of the Triton kernel is insensitive to the execution sequence. Otherwise, a deadlock may occur.\| 0: The optimization is disabled.<br>1: Enabled.\| \|
	31	+\| Running and scheduling\| TRITON_ENABLE_TASKQUEUE \| 0 or not set\| Specifies whether to enable task_queue.\| 0: The optimization is disabled.<br>1: Enabled.\| \|
	32	+\| Running and scheduling\| TRITON_ENABLE_SANITIZER \| 0 or not set\| Specifies whether to enable SANITIZER.\| 0: The optimization is disabled.<br>1: Enabled.\| \|
	33	+\| Running and scheduling\| ENABLE_PRINT_UB_BITS \| 0 or not set\| After this parameter is enabled, the current UB usage can be obtained for the inductor.\| 0: The optimization is disabled.<br>1: Enabled.\| \|
	34	+\| Others\| TRITON_BENCH_METHOD \| Not set\| When the Ascend NPU is used, change `do_bench` in `testing.py` to `do_bench_npu`. (This parameter is used when `INDUCTOR_ASCEND_AGGRESSIVE_AUTOTUNE` is set to `1`.) If this parameter is set to `default`, the original `do_bench` function is still called even if the NPU is available.\| "npu": Switch to `do_bench_npu`.\| \|
	35	+\| Others\| TRITON_REMOTE_RUN_CONFIG_PATH \| path \| Specifies the configuration path for remote running.\| Specify the path directly.\| \|

@@ -0,0 +1,12 @@
	1	+# Examples
	2	+
	3	+## Triton Examples
	4	+\|Example\|Description\|
	5	+\|--\|--\|
	6	+\| [01_vector_add_example](./01_vector_add_example.md)\| Example of simple addition\|
	7	+\| [02_fused_softmax_example](./02_fused_softmax_example.md)\| Example of Softmax fused operator\|
	8	+\| [03_layer_norm_example](./03_layer_norm_example.md)\| Example of Layer Normalization\|
	9	+\| [04_fused_attention_example](./04_fused_attention_example.md)\| Example of Flash Attention v2 fused attention algorithm\|
	10	+\| [05_matrix_multiplication_example](./05_matrix_multiplication_example.md)\| Example of efficient matrix multiplication\|
	11	+\| [06_autotune_example](./06_autotune_example.md)\| Example of using Autotune on kernels\|
	12	+\| [07_accuracy_comparison_example](./07_accuracy_comparison_example.md)\| Accuracy comparison example\|