Star725
1026
代码
代码
Issues61
Pull Requests202
Actions
项目讨论
Wiki
分析
Star725
1026
已合并
Fix:修改Ascend C编程指南中编译与运行章节 #1939
wk0911创建于 5月5日
Fix:修改Ascend C编程指南中编译与运行章节 #1939
已合并
wk0911创建于 5月5日
从wk0911/asc-devkitmaster合入到cann/asc-devkitmaster
共 26 个文件变更+813-1029
Adocs/guide/figures/aicore_compilation.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:7088c41a06d8b2bad800f18d1d8f7da692cf9905cf0a962dd75341bac958d982
3+size 100653
Adocs/guide/figures/aicore_simt_comilation.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:449048944d499ad9e16ce999b177020630d31af0d8b4cd9f4c62f0cf6393f5c5
3+size 87587
Adocs/guide/figures/asynchronous_execution.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:ee0bb20b29aefd8771897a53972a4863a41a3591ef2d94fcc06ff6ab8a723911
3+size 99318
Adocs/guide/figures/asynchronous_execution_async.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:3b475c24a937549a35c63f057769315d4ce8508fd0b858a87362de84c31debb6
3+size 102812
Mdocs/guide/index.md+8-18
@@ -36,24 +36,14 @@
36 36 
37 - [AI CPU编程](编程指南/编程模型/AI-CPU编程.md)37 - [AI CPU编程](编程指南/编程模型/AI-CPU编程.md)
38 38 
39- - [编译与运行](编程指南/编译与运行/AI-Core-SIMD编译/算子编译简介.md)39+ - [编译与运行](编程指南/编译与运行/异步执行.md)
40- - [AI Core SIMD编译](编程指南/编译与运行/AI-Core-SIMD编译/算子编译简介.md)40+ - [异步执行](编程指南/编译与运行/异步执行.md)
41- - [算子编译简介](编程指南/编译与运行/AI-Core-SIMD编译/算子编译简介.md)41+ - [算子编译](编程指南/编译与运行/算子编译/毕昇编译器.md)
42- - [通过bisheng命令行编译](编程指南/编译与运行/AI-Core-SIMD编译/通过bisheng命令行编译.md)42+ - [毕昇编译器](编程指南/编译与运行/算子编译/毕昇编译器.md)
43- - [常用的编译选项](编程指南/编译与运行/AI-Core-SIMD编译/常用的编译选项.md)43+ - [AI Core算子编译基本用法](编程指南/编译与运行/算子编译/AI-Core算子编译基本用法.md)
44- - [通过CMake编译](编程指南/编译与运行/AI-Core-SIMD编译/通过CMake编译.md)44+ - [AI CPU算子编译基本用法](编程指南/编译与运行/算子编译/AI-CPU算子编译基本用法.md)
45- - [RTC](编程指南/编译与运行/AI-Core-SIMD编译/RTC.md)45+ - [RTC运行时编译](编程指南/编译与运行/算子编译/RTC运行时编译.md)
46- - [约束说明](编程指南/编译与运行/AI-Core-SIMD编译/约束说明.md)46+ - [约束说明](编程指南/编译与运行/算子编译/约束说明.md)
47- 
48- - [AI Core SIMT编译](编程指南/编译与运行/AI-Core-SIMT编译/算子编译简介.md)
49- - [算子编译简介](编程指南/编译与运行/AI-Core-SIMT编译/算子编译简介.md)
50- - [通过bisheng命令行编译](编程指南/编译与运行/AI-Core-SIMT编译/通过bisheng命令行编译.md)
51- - [常用的编译选项](编程指南/编译与运行/AI-Core-SIMT编译/常用的编译选项.md)
52- - [通过CMake编译](编程指南/编译与运行/AI-Core-SIMT编译/通过CMake编译.md)
53- - [约束说明](编程指南/编译与运行/AI-Core-SIMT编译/约束说明.md)
54- 
55- - [AI CPU算子编译](编程指南/编译与运行/AI-CPU算子编译.md)
56- - [算子运行](编程指南/编译与运行/算子运行.md)
57 47 
58 - [语言扩展层](编程指南/语言扩展层/SIMD-BuiltIn关键字.md)48 - [语言扩展层](编程指南/语言扩展层/SIMD-BuiltIn关键字.md)
59 - [SIMD BuiltIn关键字](编程指南/语言扩展层/SIMD-BuiltIn关键字.md)49 - [SIMD BuiltIn关键字](编程指南/语言扩展层/SIMD-BuiltIn关键字.md)
Mdocs/guide/编程指南/编程模型/AI-CPU编程.md+1-2
@@ -52,7 +52,7 @@ hello_world<<<numBlocks, nullptr, stream>>>(&args, sizeof(KernelArgs));
52>- Host侧调用\_\_global\_\_ \_\_aicpu\_\_函数时必须使用<<<\>\>\>异构调用语法，输入的函数入参在入参指针的基础上需要输入从指针中读取的数据大小。52>- Host侧调用\_\_global\_\_ \_\_aicpu\_\_函数时必须使用<<<\>\>\>异构调用语法，输入的函数入参在入参指针的基础上需要输入从指针中读取的数据大小。
53>- 在Host侧使用内核调用符<<<...\>\>\>调用AI Core与AI CPU算子时不能使用同一条stream。53>- 在Host侧使用内核调用符<<<...\>\>\>调用AI Core与AI CPU算子时不能使用同一条stream。
54 54 
55-加载和运行算子时，需要使用Runtime API，完成运行时管理和配置，详细内容请参考[算子运行](../编译与运行/算子运行.md)。AI CPU算子的编译请参考[AI CPU算子编译](../编译与运行/AI-CPU算子编译.md)。55+加载和运行算子时，需要使用Runtime API，完成运行时管理和配置，详细内容请参考[算子运行](../编译与运行/异步执行.md)。AI CPU算子的编译请参考[AI CPU算子编译](../编译与运行/算子编译/AI-CPU算子编译基本用法.md)。
56 56 
57## AI CPU模板核函数<a name="section135075471718"></a>57## AI CPU模板核函数<a name="section135075471718"></a>
58 58 
@@ -82,4 +82,3 @@ template extern __global__ __aicpu__ uint32_t hello_world<KernelArgs, 4096>(void
82 82 
83>[!NOTE]说明 83>[!NOTE]说明
84>更多AI CPU API的使用方法请参考[AI CPU API](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/context/AI-CPU-API.md)。84>更多AI CPU API的使用方法请参考[AI CPU API](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/context/AI-CPU-API.md)。
85- 
Ddocs/guide/编程指南/编译与运行/AI-Core-SIMD编译/RTC.md+0-219
@@ -1,219 +0,0 @@
1-# RTC<a name="ZH-CN_TOPIC_0000002462620725"></a>
2- 
3-RTC是Ascend C运行时编译库，通过[aclrtc](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/context/RTC.md)接口，在程序运行时，将中间代码动态编译成目标机器码，提升程序运行性能。
4- 
5-运行时编译库提供以下核心接口：
6- 
7-- aclrtcCreateProg：根据输入参数（字符串形式表达的Ascend C源代码等）创建aclrtcProg程序实例。
8-- aclrtcCompileProg：编译给定的程序，支持用户自定义编译选项，比如指定NPU架构版本号：--npu-arch=dav-2201。支持的编译选项可以参考[《毕昇编译器用户指南》](https://www.hiascend.com/document/redirect/CannCommunityBiSheng)。
9-- aclrtcGetBinDataSize：获取编译后的Device侧二进制数据的大小。
10-- aclrtcGetBinData：获取编译后的Device侧二进制数据。
11-- aclrtcDestroyProg：在编译和执行过程结束后，销毁给定的程序。
12- 
13-编译完成后需要调用如下接口完成（仅列出核心接口）Kernel加载与执行。完整流程和详细接口说明请参考[《Runtime运行时API》](https://hiascend.com/document/redirect/CannCommunityRuntimeApi)中的“Kernel加载与执行”章节。
14- 
15-1. 通过aclrtBinaryLoadFromData接口解析由aclrtcGetBinData接口获取的算子二进制数据。
16-2. 获取核函数句柄并根据核函数句柄操作其参数列表，相关接口包括aclrtBinaryGetFunction（获取核函数句柄）、aclrtKernelArgsInit（初始化参数列表）、aclrtKernelArgsAppend（追加拷贝用户设置的参数值如xDevice, yDevice, zDevice）等。
17-3. 调用aclrtLaunchKernelWithConfig接口，启动对应算子的计算任务。
18- 
19-如下是一个使用aclrtc接口编译并运行Add自定义算子的完整样例：
20- 
21-```
22-#include <iostream>
23-#include <fstream>
24-#include <vector>
25-#include "acl/acl.h"
26-// 使用aclrtc接口需要包含的头文件
27-#include "acl/acl_rt_compile.h"
28- 
29-#define CHECK_ACL(x) \
30- do { \
31- aclError __ret = x; \
32- if (__ret != ACL_ERROR_NONE) { \
33- std::cerr << __FILE__ << ":" << __LINE__ << " aclError:" << __ret << std::endl; \
34- } \
35- } while (0);
36- 
37-int main(int argc, char *argv[])
38-{
39- // ----------------------------------------------------- aclrtc part -----------------------------------------------------
40- const char *src = R""""(
41-#include "kernel_operator.h"
42-constexpr int32_t TOTAL_LENGTH = 8 * 1024; // total length of data
43-constexpr int32_t USE_CORE_NUM = 8; // num of core used
44-constexpr int32_t BLOCK_LENGTH = TOTAL_LENGTH / USE_CORE_NUM; // length computed of each core
45-constexpr int32_t TILE_NUM = 8; // split data into 8 tiles for each core
46-constexpr int32_t BUFFER_NUM = 2; // tensor num for each queue
47-constexpr int32_t TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM; // separate to 2 parts, due to double buffer
48- 
49-class KernelAdd {
50-public:
51- __aicore__ inline KernelAdd() {}
52- __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z)
53- {
54- xGm.SetGlobalBuffer((__gm__ float *)x + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
55- yGm.SetGlobalBuffer((__gm__ float *)y + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
56- zGm.SetGlobalBuffer((__gm__ float *)z + BLOCK_LENGTH * AscendC::GetBlockIdx(), BLOCK_LENGTH);
57- pipe.InitBuffer(inQueueX, BUFFER_NUM, TILE_LENGTH * sizeof(float));
58- pipe.InitBuffer(inQueueY, BUFFER_NUM, TILE_LENGTH * sizeof(float));
59- pipe.InitBuffer(outQueueZ, BUFFER_NUM, TILE_LENGTH * sizeof(float));
60- }
61- __aicore__ inline void Process()
62- {
63- int32_t loopCount = TILE_NUM * BUFFER_NUM;
64- for (int32_t i = 0; i < loopCount; i++) {
65- CopyIn(i);
66- Compute(i);
67- CopyOut(i);
68- }
69- }
70- 
71-private:
72- __aicore__ inline void CopyIn(int32_t progress)
73- {
74- AscendC::LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
75- AscendC::LocalTensor<float> yLocal = inQueueY.AllocTensor<float>();
76- AscendC::DataCopy(xLocal, xGm[progress * TILE_LENGTH], TILE_LENGTH);
77- AscendC::DataCopy(yLocal, yGm[progress * TILE_LENGTH], TILE_LENGTH);
78- inQueueX.EnQue(xLocal);
79- inQueueY.EnQue(yLocal);
80- }
81- __aicore__ inline void Compute(int32_t progress)
82- {
83- AscendC::LocalTensor<float> xLocal = inQueueX.DeQue<float>();
84- AscendC::LocalTensor<float> yLocal = inQueueY.DeQue<float>();
85- AscendC::LocalTensor<float> zLocal = outQueueZ.AllocTensor<float>();
86- AscendC::Add(zLocal, xLocal, yLocal, TILE_LENGTH);
87- 
88- outQueueZ.EnQue<float>(zLocal);
89- inQueueX.FreeTensor(xLocal);
90- inQueueY.FreeTensor(yLocal);
91- }
92- __aicore__ inline void CopyOut(int32_t progress)
93- {
94- AscendC::LocalTensor<float> zLocal = outQueueZ.DeQue<float>();
95- AscendC::DataCopy(zGm[progress * TILE_LENGTH], zLocal, TILE_LENGTH);
96- outQueueZ.FreeTensor(zLocal);
97- }
98- 
99-private:
100- AscendC::TPipe pipe;
101- AscendC::TQue<AscendC::TPosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;
102- AscendC::TQue<AscendC::TPosition::VECOUT, BUFFER_NUM> outQueueZ;
103- AscendC::GlobalTensor<float> xGm;
104- AscendC::GlobalTensor<float> yGm;
105- AscendC::GlobalTensor<float> zGm;
106-};
107-extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
108-{
109- KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY);
110- KernelAdd op;
111- op.Init(x, y, z);
112- op.Process();
113-}
114-)"""";
115- // aclrtc流程，src为用户Device侧源码，通过aclrtcCreateProg来创建编译程序
116- aclrtcProg prog;
117- CHECK_ACL(aclrtcCreateProg(&prog, src, "add_custom", 0, nullptr, nullptr));
118- 
119- // aclrtc流程，传入毕昇编译器的编译选项，调用aclrtcCompileProg进行编译
120- const char *options[] = {
121- "--npu-arch=dav-2201",
122- };
123- int numOptions = sizeof(options) / sizeof(options[0]);
124- CHECK_ACL(aclrtcCompileProg(prog, numOptions, options));
125- 
126- // aclrtc流程，获取Device侧二进制内容和大小
127- size_t binDataSizeRet;
128- CHECK_ACL(aclrtcGetBinDataSize(prog, &binDataSizeRet));
129- std::vector<char> deviceELF(binDataSizeRet);
130- CHECK_ACL(aclrtcGetBinData(prog, deviceELF.data()));
131- 
132- const char *funcName = "add_custom";
133- // ----------------------------------------------------- aclrt part -----------------------------------------------------
134- uint32_t numBlocks = 8;
135- size_t inputByteSize = 8 * 1024 * sizeof(uint32_t);
136- size_t outputByteSize = 8 * 1024 * sizeof(uint32_t);
137- CHECK_ACL(aclInit(nullptr));
138- int32_t deviceId = 0;
139- CHECK_ACL(aclrtSetDevice(deviceId));
140- aclrtStream stream = nullptr;
141- CHECK_ACL(aclrtCreateStream(&stream));
142- 
143- uint8_t *xHost, *yHost, *zHost;
144- uint8_t *xDevice, *yDevice, *zDevice;
145- CHECK_ACL(aclrtMallocHost((void **)(&xHost), inputByteSize));
146- CHECK_ACL(aclrtMallocHost((void **)(&yHost), inputByteSize));
147- CHECK_ACL(aclrtMallocHost((void **)(&zHost), outputByteSize));
148- CHECK_ACL(aclrtMalloc((void **)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
149- CHECK_ACL(aclrtMalloc((void **)&yDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
150- CHECK_ACL(aclrtMalloc((void **)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
151- CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
152- CHECK_ACL(aclrtMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
153- 
154- aclrtBinHandle binHandle = nullptr;
155- aclrtBinaryLoadOptions loadOption;
156- loadOption.numOpt = 1;
157- aclrtBinaryLoadOption option;
158- option.type = ACL_RT_BINARY_LOAD_OPT_LAZY_MAGIC;
159- option.value.magic = ACL_RT_BINARY_MAGIC_ELF_VECTOR_CORE; // 设置magic值，表示算子在Vector Core上执行
160- loadOption.options = &option;
161- CHECK_ACL(aclrtBinaryLoadFromData(deviceELF.data(), binDataSizeRet, &loadOption, &binHandle));
162- aclrtFuncHandle funcHandle = nullptr;
163- 
164- CHECK_ACL(aclrtBinaryGetFunction(binHandle, funcName, &funcHandle));
165- 
166- aclrtArgsHandle argsHandle = nullptr;
167- aclrtParamHandle paramHandle = nullptr;
168- CHECK_ACL(aclrtKernelArgsInit(funcHandle, &argsHandle));
169- CHECK_ACL(aclrtKernelArgsAppend(argsHandle, (void **)&xDevice, sizeof(uintptr_t), &paramHandle));
170- CHECK_ACL(aclrtKernelArgsAppend(argsHandle, (void **)&yDevice, sizeof(uintptr_t), &paramHandle));
171- CHECK_ACL(aclrtKernelArgsAppend(argsHandle, (void **)&zDevice, sizeof(uintptr_t), &paramHandle));
172- CHECK_ACL(aclrtKernelArgsFinalize(argsHandle));
173- // 核函数入口
174- CHECK_ACL(aclrtLaunchKernelWithConfig(funcHandle, numBlocks, stream, nullptr, argsHandle, nullptr));
175- 
176- CHECK_ACL(aclrtSynchronizeStream(stream));
177- CHECK_ACL(aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST));
178- 
179- // 获取日志大小并得到日志字符串
180- size_t logSize;
181- CHECK_ACL(aclrtcGetCompileLogSize(prog, &logSize));
182- char* log = (char*)malloc(logSize);
183- CHECK_ACL(aclrtcGetCompileLog(prog, log));
184- // 将日志字符串存到文件中
185- /*
186- std::ofstream logFile("compile.log");
187- if (logFile.is_open()) {
188- logFile << log << std::endl;
189- logFile.close();
190- std::cout << "already write to compile.log!" << std::endl;
191- }
192- */
193- free(log);
194- CHECK_ACL(aclrtBinaryUnLoad(binHandle));
195- CHECK_ACL(aclrtFree(xDevice));
196- CHECK_ACL(aclrtFree(yDevice));
197- CHECK_ACL(aclrtFree(zDevice));
198- CHECK_ACL(aclrtFreeHost(xHost));
199- CHECK_ACL(aclrtFreeHost(yHost));
200- CHECK_ACL(aclrtFreeHost(zHost));
201- CHECK_ACL(aclrtDestroyStream(stream));
202- CHECK_ACL(aclrtResetDevice(deviceId));
203- CHECK_ACL(aclFinalize());
204- 
205- // 编译和运行均已结束，销毁程序
206- CHECK_ACL(aclrtcDestroyProg(&prog));
207- 
208- return 0;
209-}
210-```
211- 
212-编译命令如下，编译时需要设置-I路径为$\{INSTALL\_DIR\}/include，用于找到aclrtc相关头文件，并需要链接alc\_rtc动态库。
213- 
214-```
215-g++ add_custom.cpp -I${INSTALL_DIR}/include -L${INSTALL_DIR}/lib64 -lascendcl -lacl_rtc -o main
216-```
217- 
218-$\{INSTALL\_DIR\}请替换为CANN软件安装后文件存储路径。以root用户安装为例，安装后文件默认存储路径为：/usr/local/Ascend/cann。
219- 
Ddocs/guide/编程指南/编译与运行/AI-Core-SIMD编译/常用的编译选项.md+0-65
@@ -1,65 +0,0 @@
1-# 常用的编译选项
2- 
3-## 常用的编译选项
4- 
5-常用的编译选项说明如下，全量的编译选项请参考[《毕昇编译器用户指南》](https://www.hiascend.com/document/redirect/CannCommunityBiSheng)。
6- 
7-| 选项 | 是否必需 | 说明 |
8-| --- | --- | --- |
9-| -help | 否 | 查看帮助。 |
10-| --npu-arch | 是 | 编译时指定的AI处理器架构，取值为dav-，其中为NPU架构版本号，各产品型号对应的架构版本号请通过[对应关系表](SIMD-BuiltIn关键字.md#table65291052154114)进行查询。 |
11-| --npu-soc | 否 | 编译时指定的AI处理器型号，npu-soc和npu-arch同时配置时，优先使能npu-arch。<br>AI处理器的型号请通过如下方式获取：<br>- 针对<term>Atlas 推理系列产品</term>、<term>Atlas A2 训练系列产品</term>/<term>Atlas A2 推理系列产品</term>：在安装AI处理器的服务器执行npu-smi info命令进行查询，获取Name信息。实际配置值为AscendName，例如Name取值为xxxyy，实际配置值为Ascendxxxyy。<br><br>- 针对<term>Atlas A3 训练系列产品</term>/<term>Atlas A3 推理系列产品</term>，在安装AI处理器的服务器执行npu-smi info -t board -i id -c chip_id命令进行查询，获取Chip Name和NPU Name信息，实际配置值为Chip Name_NPU Name。例如Chip Name取值为Ascendxxx，NPU Name取值为1234，实际配置值为Ascendxxx_1234。其中：<br>id：设备id，通过npu-smi info -l命令查出的NPU ID即为设备id。<br> chip_id：芯片id，通过npu-smi info -m命令查出的Chip ID即为芯片id。<br><br>- 针对Ascend 950PR/Ascend 950DT，在安装AI处理器的服务器执行npu-smi info -t board -i id命令进行查询，获取Chip Name和NPU Name信息，实际配置值为Chip Name_NPU Name。例如Chip Name取值为Ascendxxx，NPU Name取值为1234，实际配置值为Ascendxxx_1234。<br> 其中：id为设备id，通过npu-smi info -l命令查出的NPU ID即为设备id。|
12-| -x | 否 | 指定编译语言。<br>指定为asc时表示Ascend C编程语言。 |
13-| -o | 否 | 指定输出文件的名称和位置。 |
14-| -c | 否 | 编译生成目标文件。 |
15-| -shared，--shared | 否 | 编译生成动态链接库。 |
16-| -lib，--cce-build-static-lib | 否 | 编译生成静态链接库。编译器会将Device侧的代码进行编译链接，生成Device侧二进制文件，随后将该文件作为Host侧编译的输入进行编译，最后链接生成静态链接库。 |
17-| -g | 否 | 编译时增加调试信息。 |
18-| --sanitizer | 否 | 编译时增加代码正确性校验信息。使用sanitizer选项时，需要同步添加-g选项，且不能在-O0场景下使用。<br>注意，启用该选项后GlobalTensor默认使能L2 Cache，无法通过AscendC::SetL2CacheHint接口设置不使能L2 Cache的模式。 |
19-| -fPIC | 否 | 告知编译器产生位置无关代码。 |
20-| -O | 否 | 用于指定编译器的优化级别，当前支持-O3，-O2，-O0。 |
21-| --run-mode=sim | 否 | sim模式：链接时用户添加仿真模式对应的实现库，实现代码在仿真模式下运行，可以查看仿真相关日志，方便用户性能调试。 |
22- 
23-## 内置编译宏开关
24- 
25-内置编译宏开关列表如下：
26- 
27-- ASCENDC\_DUMP用于控制Dump开关，默认开关打开，开发者调用printf/DumpTensor/assert后会有信息打印（需要注意直调工程的kernel文件内存在host函数，如果在host函数内调用了printf接口，也会触发kernel内的printf相关初始化动作，进而影响kernel的执行性能\)；设置为0后，表示开关关闭。示例如下：
28- 
29- ```
30- // 关闭所有算子的printf打印功能
31- ascendc_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE
32- ASCENDC_DUMP=0
33- )
34- ```
35- 
36-- ASCENDC\_DEBUG用于控制Ascend C API的调测开关，默认开关关闭；增加该编译宏后，表示开关打开，此时接口内部的assert校验生效，校验不通过会有assert日志打屏。开启该功能会对算子实际运行的性能带来一定影响，通常在调测阶段使用。示例如下：
37- 
38- ```
39- ascendc_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE
40- ASCENDC_DEBUG
41- )
42- ```
43- 
44- 当前ASCENDC\_DEBUG功能支持的产品型号为：
45- 
46- Atlas 推理系列产品
47- 
48- Atlas A2 训练系列产品/Atlas A2 推理系列产品
49- 
50-- ENABLE\_CV\_COMM\_VIA\_SSBUF用于控制是否使用SSBuffer以及UB到L1 Buffer的硬通道，在涉及CV通信（AIC和AIV）或使用数据搬运API时需关注此选项。开启该选项可以提高相关API的性能或拓展使用更多功能。默认开关关闭；设置为true后，表示开关打开。示例如下：
51- 
52- ```
53- ascendc_compile_definitions(ascendc_kernels_${RUN_MODE} PRIVATE
54- ENABLE_CV_COMM_VIA_SSBUF=true
55- )
56- ```
57- 
58- 仅在Ascend 950PR/Ascend 950DT支持该选项。
59- 
60- - 从其它硬件平台移植到此平台的算子，开关默认关闭以保持兼容性。
61- - 在该平台新开发的算子，以下场景需要打开：使用矩阵计算Matmul高阶API，且使用SetTensorScaleA等接口，这些接口属于Ascend 950PR/Ascend 950DT新增的功能，其内部实现使用了SSBuffer；使用DataCopy接口从UB拷贝数据到L1 Buffer。
62- 
63-- NO\_OVERLAP\_IN\_MULTI\_REPEAT
64- 
65- 该编译选项用于在没有地址重叠的情况下移除不必要的内存同步指令，以提升性能。针对Ascend 950PR/Ascend 950DT，使用基础API的高维切分计算API时，默认会插入内存同步指令以确保在地址重叠等复杂场景下的数据正确性，但这些同步指令会带来性能开销。在追求极致性能的场景下，如果您可以确定代码在任何情况下都不会发生内存重叠，可以使用此选项。
Ddocs/guide/编程指南/编译与运行/AI-Core-SIMD编译/算子编译简介.md+0-12
@@ -1,12 +0,0 @@
1-# 算子编译简介<a name="ZH-CN_TOPIC_0000002457558450"></a>
2- 
3-本章节介绍的算子编译方法支持开发者通过bisheng命令行和CMake进行手动配置编译选项，或编写CMake脚本来实现编译。开发者可以将Host侧main.cpp和Device侧Kernel核函数置于同一实现文件中，以实现异构编译。
4- 
5-- 目前，该编译方法仅支持如下型号：
6- - Ascend 950PR/Ascend 950DT
7- - Atlas A3 训练系列产品/Atlas A3 推理系列产品
8- - Atlas A2 训练系列产品/Atlas A2 推理系列产品
9- - Atlas 推理系列产品
10- 
11-- 异构编译场景中的编程相关约束请参考[约束说明](约束说明.md)。
12- 
Ddocs/guide/编程指南/编译与运行/AI-Core-SIMD编译/通过CMake编译.md+0-162
@@ -1,162 +0,0 @@
1-# 通过CMake编译<a name="ZH-CN_TOPIC_0000002428982142"></a>
2- 
3-项目中可以使用CMake来更简便地使用毕昇编译器编译Ascend C算子，生成可执行文件、动态库、静态库或二进制文件。
4- 
5-以下是CMake脚本的示例及其核心步骤说明：
6- 
7-```
8-# 1、find_package(ASC)是CMake中用于查找和配置Ascend C编译工具链的命令
9-find_package(ASC)
10- 
11-# 2、指定项目支持的语言包括ASC和CXX，ASC表示支持使用毕昇编译器对Ascend C编程语言进行编译
12-project(kernel_samples LANGUAGES ASC CXX)
13- 
14-# 3、使用CMake接口编译可执行文件、动态库、静态库、二进制文件
15-add_executable(demo
16- add_custom.asc
17-)
18-#.....
19-target_compile_options(demo PRIVATE
20- # --npu-arch用于指定NPU的架构版本，dav-后为架构版本号，各产品型号对应的架构版本号请通过[对应关系表](../../语言扩展层/SIMD-BuiltIn关键字.md#table65291052154114)进行查询。
21- # <COMPILE_LANGUAGE:ASC>:表明该编译选项仅对语言ASC生效
22- $<$<COMPILE_LANGUAGE:ASC>: --npu-arch=dav-2201>
23-)
24-```
25- 
26-以下是动态库、静态库编译示例，同时展示如何将源文件切换为用语言ASC编译：
27- 
28-- 编译.cpp文件生成动态库
29- 
30- ```
31- # 将.cpp文件置为ASC属性，启用Ascend C语言进行编译
32- set_source_files_properties(
33- add_custom_base.cpp
34- sub_custom_base.cpp
35- PROPERTIES LANGUAGE ASC
36- )
37-
38- add_library(kernel_lib SHARED
39- add_custom_base.cpp
40- sub_custom_base.cpp
41- )
42-
43- target_compile_options(kernel_lib PRIVATE
44- $<$<COMPILE_LANGUAGE:ASC>: --npu-arch=dav-2201>
45- )
46-
47- add_executable(demo
48- main.cpp
49- )
50- target_link_libraries(demo PRIVATE
51- kernel_lib
52- )
53- ```
54- 
55-- 编译.asc文件生成静态库
56- 
57- ```
58- # .asc文件会默认启用Ascend C语言进行编译，不需要通过set_source_files_properties进行设置
59- add_library(kernel_lib STATIC
60- add_custom_base.asc
61- sub_custom_base.asc
62- )
63-
64- target_compile_options(kernel_lib PRIVATE
65- $<$<COMPILE_LANGUAGE:ASC>: --npu-arch=dav-2201>
66- )
67-
68- add_executable(demo
69- main.cpp
70- )
71- target_link_libraries(demo PRIVATE
72- kernel_lib
73- )
74- ```
75- 
76-下文列出了使用CMake编译时常用的链接库、以及默认链接库。
77- 
78-**表 1** 常用的链接库（在使用高阶API时，必须链接以下库，因为这些库是高阶API功能所依赖的。在其他场景下，可以根据具体需求选择是否链接这些库。）
79- 
80-<a name="table1632724817328"></a>
81-<table><thead align="left"><tr id="row1132712485321"><th class="cellrowborder" valign="top" width="42.19%" id="mcps1.2.3.1.1"><p id="p377063619281"><a name="p377063619281"></a><a name="p377063619281"></a>使用场景</p>
82-</th>
83-<th class="cellrowborder" valign="top" width="57.809999999999995%" id="mcps1.2.3.1.2"><p id="p5423124282814"><a name="p5423124282814"></a><a name="p5423124282814"></a>名称</p>
84-</th>
85-</tr>
86-</thead>
87-<tbody><tr id="row131961149143112"><td class="cellrowborder" rowspan="3" valign="top" width="42.19%" headers="mcps1.2.3.1.1 "><p id="p477013682820"><a name="p477013682820"></a><a name="p477013682820"></a>使用高阶API相关的Tiling接口时需要同时链接。</p>
88-</td>
89-<td class="cellrowborder" valign="top" width="57.809999999999995%" headers="mcps1.2.3.1.2 "><p id="p94248427280"><a name="p94248427280"></a><a name="p94248427280"></a>libtiling_api.a</p>
90-</td>
91-</tr>
92-<tr id="row957113206355"><td class="cellrowborder" valign="top" headers="mcps1.2.3.1.1 "><p id="p1942416423283"><a name="p1942416423283"></a><a name="p1942416423283"></a>libregister.so</p>
93-</td>
94-</tr>
95-<tr id="row1652295021916"><td class="cellrowborder" valign="top" headers="mcps1.2.3.1.1 "><p id="p134241942162817"><a name="p134241942162817"></a><a name="p134241942162817"></a>libgraph_base.so</p>
96-</td>
97-</tr>
98-<tr id="row025115214213"><td class="cellrowborder" valign="top" width="42.19%" headers="mcps1.2.3.1.1 "><p id="p1777015364285"><a name="p1777015364285"></a><a name="p1777015364285"></a>使用PlatformAscendC相关硬件平台信息接口时需要链接。</p>
99-</td>
100-<td class="cellrowborder" valign="top" width="57.809999999999995%" headers="mcps1.2.3.1.2 "><p id="p1642415421284"><a name="p1642415421284"></a><a name="p1642415421284"></a>libplatform.so</p>
101-</td>
102-</tr>
103-</tbody>
104-</table>
105- 
106-**表 2** 默认链接库
107- 
108-<a name="table201231542115513"></a>
109-<table><thead align="left"><tr id="row171231542205510"><th class="cellrowborder" valign="top" width="23.98%" id="mcps1.2.3.1.1"><p id="p11123114295513"><a name="p11123114295513"></a><a name="p11123114295513"></a>名称</p>
110-</th>
111-<th class="cellrowborder" valign="top" width="76.02%" id="mcps1.2.3.1.2"><p id="p1412374225512"><a name="p1412374225512"></a><a name="p1412374225512"></a>作用描述</p>
112-</th>
113-</tr>
114-</thead>
115-<tbody><tr id="row5123842135514"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p1212364212559"><a name="p1212364212559"></a><a name="p1212364212559"></a>libascendc_runtime.a</p>
116-</td>
117-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p1112394218551"><a name="p1112394218551"></a><a name="p1112394218551"></a>Ascend C算子参数等组装库。</p>
118-</td>
119-</tr>
120-<tr id="row612324285519"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p01231423552"><a name="p01231423552"></a><a name="p01231423552"></a>libruntime.so</p>
121-</td>
122-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p8123164255511"><a name="p8123164255511"></a><a name="p8123164255511"></a>Runtime运行库。</p>
123-</td>
124-</tr>
125-<tr id="row1612374285512"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p2012315425551"><a name="p2012315425551"></a><a name="p2012315425551"></a>libprofapi.so</p>
126-</td>
127-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p12123164265514"><a name="p12123164265514"></a><a name="p12123164265514"></a>Ascend C算子运行性能数据采集库。</p>
128-</td>
129-</tr>
130-<tr id="row10123134212552"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p812374235515"><a name="p812374235515"></a><a name="p812374235515"></a>libunified_dlog.so</p>
131-</td>
132-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p412314426554"><a name="p412314426554"></a><a name="p412314426554"></a>CANN日志收集库。</p>
133-</td>
134-</tr>
135-<tr id="row1012384210552"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p15123104219559"><a name="p15123104219559"></a><a name="p15123104219559"></a>libmmpa.so</p>
136-</td>
137-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p13123242135519"><a name="p13123242135519"></a><a name="p13123242135519"></a>CANN系统接口库。</p>
138-</td>
139-</tr>
140-<tr id="row17124154245516"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p612484265518"><a name="p612484265518"></a><a name="p612484265518"></a>libascend_dump.so</p>
141-</td>
142-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p101241842175512"><a name="p101241842175512"></a><a name="p101241842175512"></a>CANN维测信息库。</p>
143-</td>
144-</tr>
145-<tr id="row6124164213551"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p111246426558"><a name="p111246426558"></a><a name="p111246426558"></a>libc_sec.so</p>
146-</td>
147-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p31241442185512"><a name="p31241442185512"></a><a name="p31241442185512"></a>CANN安全函数库。</p>
148-</td>
149-</tr>
150-<tr id="row171241342175514"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p6124124218556"><a name="p6124124218556"></a><a name="p6124124218556"></a>liberror_manager.so</p>
151-</td>
152-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p61248424557"><a name="p61248424557"></a><a name="p61248424557"></a>CANN错误信息管理库。</p>
153-</td>
154-</tr>
155-<tr id="row512404213550"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p151243425553"><a name="p151243425553"></a><a name="p151243425553"></a>libascendcl.so</p>
156-</td>
157-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p1012424213555"><a name="p1012424213555"></a><a name="p1012424213555"></a>acl相关接口库。</p>
158-</td>
159-</tr>
160-</tbody>
161-</table>
162- 
Ddocs/guide/编程指南/编译与运行/AI-Core-SIMD编译/通过bisheng命令行编译.md+0-147
@@ -1,147 +0,0 @@
1-# 通过bisheng命令行编译<a name="ZH-CN_TOPIC_0000002429141954"></a>
2- 
3-毕昇编译器是一款专为AI处理器设计的编译器，支持异构编程扩展，可以将用户编写的昇腾算子代码编译成二进制可执行文件和动态库等形式。毕昇编译器的可执行程序命名为bisheng，支持x86、aarch64等主机系统，并且原生支持设备侧AI Core架构指令集编译。通过使用毕昇编译器，用户可以更加高效地进行针对昇腾AI处理器的编程和开发工作。
4- 
5-## 入门示例<a name="section17595637204517"></a>
6- 
7-以下是一个使用毕昇编译器编译静态Shape的add\_custom算子入门示例。该示例展示了如何编写源文件add\_custom.asc以及具体的编译命令。通过这个示例，您可以了解如何使用毕昇编译器进行算子编译。完整样例请参考[LINK](https://gitcode.com/cann/asc-devkit/tree/master/examples/01_simd_cpp_api/00_introduction/01_vector/basic_api_tque_add)。
8- 
9-1. 包含头文件。
10- 
11- 在编写算子源文件时，需要包含必要的头文件。
12- 
13- ```
14- // 头文件
15- #include "acl/acl.h"
16- #include "kernel_operator.h"
17- ```
18- 
19-2. 核函数实现。
20- 
21- - 核函数支持模板。
22- - 核函数入参支持传入用户自定义的结构体，比如示例中用户自定义的AddCustomTilingData结构体。
23- 
24- ```
25- // 用户自定义的TilingData结构体
26- struct AddCustomTilingData {
27- uint32_t totalLength;
28- uint32_t tileNum;
29- };
30-
31- // Kernel核心实现逻辑，包括搬运，计算等
32- class KernelAdd {
33- public:
34- __aicore__ inline KernelAdd() {}
35- // ...
36-
37- };
38-
39- __global__ __vector__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, AddCustomTilingData tiling) // 该算子执行时仅启动AI Core上的Vector核
40- {
41-
42- KernelAdd op;
43- op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
44- op.Process();
45- }
46- ```
47- 
48-3. Host侧调用函数逻辑，包括内存申请和释放，初始化和去初始化，内核调用符调用核函数等。
49- 
50- ```
51- // Host侧应用程序需要包含的头文件
52- #include "acl/acl.h"
53- // Kernel侧需要包含的头文件
54- #include "kernel_operator.h"
55- // 核函数开发部分
56- ...
57-
58- __global__ __vector__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z, AddCustomTilingData tiling)
59- {
60-
61- KernelAdd op;
62- op.Init(x, y, z, tiling.totalLength, tiling.tileNum);
63- op.Process();
64- }
65-
66- // 通过<<<...>>>内核调用符调用算子
67- std::vector<float> kernel_add(std::vector<float> &x, std::vector<float> &y)
68- {
69- ...
70- }
71-
72-
73- // 计算结果比对
74- uint32_t VerifyResult(std::vector<float> &output, std::vector<float> &golden)
75- {
76- ...
77- }
78-
79- // 算子验证主程序
80- int32_t main(int32_t argc, char *argv[])
81- {
82- constexpr uint32_t totalLength = 8 * 2048;
83- constexpr float valueX = 1.2f;
84- constexpr float valueY = 2.3f;
85- std::vector<float> x(totalLength, valueX);
86- std::vector<float> y(totalLength, valueY);
87-
88- std::vector<float> output = kernel_add(x, y);
89-
90- std::vector<float> golden(totalLength, valueX + valueY);
91- return VerifyResult(output, golden);
92- }
93- ```
94- 
95-4. 采用如下的编译命令进行编译。
96- 
97- - -o demo：指定输出文件名为demo。
98- - --npu-arch=dav-2201：指定NPU的架构版本为dav-2201。dav-后为NPU架构版本号，各产品型号对应的架构版本号请通过[对应关系表](../../语言扩展层/SIMD-BuiltIn关键字.md#table65291052154114)进行查询。
99- 
100- ```
101- bisheng add_custom.asc -o demo --npu-arch=dav-2201
102- ```
103- 
104-5. 执行可执行文件。
105- 
106- ```
107- ./demo
108- ```
109- 
110-## 程序的编译与执行<a name="section313754223513"></a>
111- 
112-通过毕昇编译器可以将算子源文件（以.asc为后缀）编译为当前平台的可执行文件或算子动态库，静态库。此外，也支持编译以.cpp/.c等为后缀的C++/C源文件，但需要增加-x asc编译选项。
113- 
114-- 编译生成可执行文件
115- 
116- ```
117- # 1.编译hello_world.cpp为当前平台可执行文件
118- # bisheng [算子源文件] -o [输出产物名称] --npu-arch=[NPU架构版本号]，常见参数顺序与g++保持一致。
119- bisheng -x asc add_custom.cpp -o add_custom --npu-arch=dav-xxxx
120- ```
121- 
122- 生成的可执行文件可通过如下方式执行：
123- 
124- ```
125- ./add_custom
126- ```
127- 
128-- 编译生成算子动态库
129- 
130- ```
131- # 2.编译add_custom_base.cpp生成算子动态库
132- # bisheng -shared [算子源文件] -o [输出产物名称] --npu-arch=[NPU架构版本号]
133- # 动态库
134- bisheng -shared -x asc add_custom_base.cpp -o libadd.so --npu-arch=dav-xxxx
135- ```
136- 
137-- 编译生成算子静态库
138- 
139- ```
140- # 3.编译add_custom_base.cpp生成算子静态库
141- bisheng -lib [算子源文件] -o [输出产物名称] --npu-arch=[NPU架构版本号]
142- # 静态库
143- bisheng -lib -x asc add_custom_base.cpp -o libadd.a --npu-arch=dav-xxxx
144- ```
145- 
146-在命令行编译场景下，可以按需链接需要的库文件，常见的库文件请参考[常用的链接库](通过CMake编译.md#table1632724817328)。编译时会默认链接[表2](通过CMake编译.md#table201231542115513)中列出的库文件。注意如下例外场景：在使用g++链接asc代码编译生成的静态库时，需要手动链接默认链接库。
147- 
Ddocs/guide/编程指南/编译与运行/AI-Core-SIMT编译/常用的编译选项.md+0-12
@@ -1,12 +0,0 @@
1-# 常用的编译选项
2- 
3-常用的编译选项说明如下，全量的编译选项请参考[《毕昇编译器用户指南》](https://www.hiascend.com/document/redirect/CannCommunityBiSheng)。
4- 
5-| 选项 | 是否必需 | 说明 |
6-| --- | --- | --- |
7-| -help | 否 | 查看帮助。 |
8-| --npu-arch | 是 | 编译时指定的AI处理器架构，取值为dav-，其中为NPU架构版本号，各产品型号对应的架构版本号请通过[对应关系表](../../语言扩展层/SIMD-BuiltIn关键字.md#table65291052154114)进行查询。 |
9-| -x | 否 | 指定编译语言。<br>指定为asc时表示Ascend C编程语言。 |
10-| -o | 否 | 指定输出文件的名称和位置。 |
11-| -c | 否 | 编译生成目标文件。 |
12-| --enable-simt | 是 | 指定SIMT编程场景。 |
Ddocs/guide/编程指南/编译与运行/AI-Core-SIMT编译/算子编译简介.md+0-9
@@ -1,9 +0,0 @@
1-# 算子编译简介<a name="ZH-CN_TOPIC_0000002561896183"></a>
2- 
3-本章节介绍的SIMT编程场景算子编译方法，支持开发者通过bisheng命令行或编写CMake脚本来实现编译。开发者可以将Host侧调用代码和Device侧Kernel核函数置于同一实现文件中，以实现异构编译。
4- 
5-- 目前，该编译方法仅支持如下型号：
6- - Ascend 950PR/Ascend 950DT
7- 
8-- 异构编译场景中的编程相关约束请参考[约束说明](约束说明.md)。
9- 
Ddocs/guide/编程指南/编译与运行/AI-Core-SIMT编译/约束说明.md+0-5
@@ -1,5 +0,0 @@
1-# 约束说明<a name="ZH-CN_TOPIC_0000002531172022"></a>
2- 
3-- bfloat16\_t等数据类型在Host侧不支持，使用这些数据类型时，Host和Device不能写在同一个实现文件里。Host侧不支持的数据类型如下：
4- 
5- Ascend 950PR/Ascend 950DT：bfloat16\_t、hifloat8\_t、fp8\_e5m2\_t、fp8\_e4m3fn\_t、fp8\_e8m0\_t、fp4x2\_e2m1\_t、fp4x2\_e1m2\_t、int4x2\_t。
Ddocs/guide/编程指南/编译与运行/AI-Core-SIMT编译/通过CMake编译.md+0-88
@@ -1,88 +0,0 @@
1-# 通过CMake编译<a name="ZH-CN_TOPIC_0000002531056204"></a>
2- 
3-项目中可以使用CMake来更简便地使用毕昇编译器编译Ascend C SIMT算子，生成可执行文件。
4- 
5-以下是CMake脚本的示例及其核心步骤说明：
6- 
7-```
8-cmake_minimum_required(VERSION 3.16)
9- 
10-# 1、find_package(ASC)是CMake中用于查找和配置Ascend C编译工具链的命令
11-find_package(ASC)
12- 
13-# 2、指定项目支持的语言包括ASC和CXX，ASC表示支持使用毕昇编译器对Ascend C编程语言进行编译
14-project(kernel_samples LANGUAGES ASC CXX)
15- 
16-# 3、使用CMake接口编译可执行文件
17-add_executable(demo
18- add_custom.asc
19-)
20-#.....
21-target_compile_options(demo PRIVATE
22- # --npu-arch用于指定NPU的架构版本，dav-后为架构版本号
23- # <COMPILE_LANGUAGE:ASC>:表明该编译选项仅对语言ASC生效
24- $<$<COMPILE_LANGUAGE:ASC>: --npu-arch=dav-3510>
25- # 开启SIMT编程模型的编译功能
26- --enable-simt
27-)
28-```
29- 
30-下文列出了使用CMake编译时默认链接库。
31- 
32-**表 1** 默认链接库
33- 
34-<a name="table201231542115513"></a>
35-<table><thead align="left"><tr id="row171231542205510"><th class="cellrowborder" valign="top" width="23.98%" id="mcps1.2.3.1.1"><p id="p11123114295513"><a name="p11123114295513"></a><a name="p11123114295513"></a>名称</p>
36-</th>
37-<th class="cellrowborder" valign="top" width="76.02%" id="mcps1.2.3.1.2"><p id="p1412374225512"><a name="p1412374225512"></a><a name="p1412374225512"></a>作用描述</p>
38-</th>
39-</tr>
40-</thead>
41-<tbody><tr id="row5123842135514"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p1212364212559"><a name="p1212364212559"></a><a name="p1212364212559"></a>libascendc_runtime.a</p>
42-</td>
43-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p1112394218551"><a name="p1112394218551"></a><a name="p1112394218551"></a>Ascend C算子参数等组装库。</p>
44-</td>
45-</tr>
46-<tr id="row612324285519"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p01231423552"><a name="p01231423552"></a><a name="p01231423552"></a>libruntime.so</p>
47-</td>
48-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p8123164255511"><a name="p8123164255511"></a><a name="p8123164255511"></a>Runtime运行库。</p>
49-</td>
50-</tr>
51-<tr id="row1612374285512"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p2012315425551"><a name="p2012315425551"></a><a name="p2012315425551"></a>libprofapi.so</p>
52-</td>
53-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p12123164265514"><a name="p12123164265514"></a><a name="p12123164265514"></a>Ascend C算子运行性能数据采集库。</p>
54-</td>
55-</tr>
56-<tr id="row10123134212552"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p812374235515"><a name="p812374235515"></a><a name="p812374235515"></a>libunified_dlog.so</p>
57-</td>
58-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p412314426554"><a name="p412314426554"></a><a name="p412314426554"></a>CANN日志收集库。</p>
59-</td>
60-</tr>
61-<tr id="row1012384210552"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p15123104219559"><a name="p15123104219559"></a><a name="p15123104219559"></a>libmmpa.so</p>
62-</td>
63-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p13123242135519"><a name="p13123242135519"></a><a name="p13123242135519"></a>CANN系统接口库。</p>
64-</td>
65-</tr>
66-<tr id="row17124154245516"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p612484265518"><a name="p612484265518"></a><a name="p612484265518"></a>libascend_dump.so</p>
67-</td>
68-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p101241842175512"><a name="p101241842175512"></a><a name="p101241842175512"></a>CANN维测信息库。</p>
69-</td>
70-</tr>
71-<tr id="row6124164213551"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p111246426558"><a name="p111246426558"></a><a name="p111246426558"></a>libc_sec.so</p>
72-</td>
73-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p31241442185512"><a name="p31241442185512"></a><a name="p31241442185512"></a>CANN安全函数库。</p>
74-</td>
75-</tr>
76-<tr id="row171241342175514"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p6124124218556"><a name="p6124124218556"></a><a name="p6124124218556"></a>liberror_manager.so</p>
77-</td>
78-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p61248424557"><a name="p61248424557"></a><a name="p61248424557"></a>CANN错误信息管理库。</p>
79-</td>
80-</tr>
81-<tr id="row512404213550"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p151243425553"><a name="p151243425553"></a><a name="p151243425553"></a>libascendcl.so</p>
82-</td>
83-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p1012424213555"><a name="p1012424213555"></a><a name="p1012424213555"></a>acl相关接口库。</p>
84-</td>
85-</tr>
86-</tbody>
87-</table>
88- 
Ddocs/guide/编程指南/编译与运行/AI-Core-SIMT编译/通过bisheng命令行编译.md+0-121
@@ -1,121 +0,0 @@
1-# 通过bisheng命令行编译<a name="ZH-CN_TOPIC_0000002561936205"></a>
2- 
3-毕昇编译器是一款专为AI处理器设计的编译器，支持异构编程扩展，可以将用户编写的昇腾算子代码编译成二进制可执行文件和动态库等形式。毕昇编译器的可执行程序命名为bisheng，支持x86、aarch64等主机系统，并且原生支持设备侧AI Core架构指令集编译。通过使用毕昇编译器，用户可以更加高效地进行针对昇腾AI处理器的编程和开发工作。
4- 
5-## 入门示例<a name="section727365944014"></a>
6- 
7-以下是一个使用毕昇编译器编译的SIMT编程实现的Add算子入门示例。该示例展示了如何编写源文件add.asc以及具体的编译命令。通过这个示例，您可以了解如何使用毕昇编译器进行SIMT算子编译。完整样例请参考[LINK](https://gitcode.com/cann/asc-devkit/tree/master/examples/03_simt_api/00_introduction/01_add)。
8- 
9-1. 包含头文件。
10- 
11- 在编写算子源文件时，需要包含必要的头文件。
12- 
13- ```
14- // 头文件
15- #include "acl/acl.h" // 调用核函数相关接口头文件
16- #include "asc_simt.h" // 核函数内调用SIMT API接口的头文件
17- ```
18- 
19-2. 核函数实现。
20- 
21- 核函数入参当前仅支持基础数据类型及其指针类型，核函数具体语法及约束说明参见[核函数](../../编程模型/AI-Core-SIMD编程/核函数.md)。
22- 
23- ```
24- __global__ void add_custom(float* x, float* y, float* z, uint64_t total_length)
25- {
26- // Calculate global thread ID
27- int32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
28- // Maps to the row index of output tensor
29- if (idx >= total_length) {
30- return;
31- }
32- z[idx] = x[idx] + y[idx];
33- }
34- ```
35- 
36-3. Host侧调用函数逻辑，包括内存申请和释放，初始化和去初始化，内核调用符调用核函数等。
37- 
38- ```
39- // Host侧应用程序需要包含的头文件
40- #include "acl/acl.h"
41-
42- // 核函数实现
43- __global__ void add_custom(float* x, float* y, float* z, uint64_t total_length)
44- {
45- ...
46- }
47-
48- // 通过<<<...>>>内核调用符调用算子
49- std::vector<float> add(std::vector<float>& x, std::vector<float>& y)
50- {
51- ...
52- // Call kernel funtion with <<<...>>>
53- add_custom<<<...>>>(...));
54- ...
55- }
56-
57-
58- // 计算结果比对
59- uint32_t verify_result(std::vector<float>& output, std::vector<float>& golden)
60- {
61- ...
62- }
63-
64- // 算子验证主程序
65- int32_t main(int32_t argc, char *argv[])
66- {
67- constexpr uint32_t in_shape = 48 * 256;
68- std::vector<float> x(in_shape);
69- for (uint32_t i = 0; i < in_shape; i++) {
70- x[i] = i * 1.1f;
71- }
72- std::vector<float> y(in_shape);
73- for (uint32_t i = 0; i < in_shape; i++) {
74- y[i] = i + 3.4f;
75- }
76- std::vector<float> golden(in_shape);
77- for (uint32_t i = 0; i < in_shape; i++) {
78- golden[i] = x[i] + y[i];
79- }
80- std::vector<float> output = add(x, y);
81- return verify_result(output, golden);
82- }
83- ```
84- 
85-4. 采用如下的编译命令进行编译。
86- 
87- ```
88- bisheng -x asc add.asc -o demo --npu-arch=dav-3510 --enable-simt
89- ```
90- 
91- - -x asc：-x指定编译语言，asc表示编程语言为Ascend C。
92- - -o demo：指定输出文件名为demo。
93- - --npu-arch=dav-3510：指定NPU的架构版本为dav-3510。dav-后为NPU架构版本号，各产品型号对应的架构版本号请通过[对应关系表](../../语言扩展层/SIMD-BuiltIn关键字.md#table65291052154114)进行查询。
94- - --enable-simt： SIMT编译的固定选项。
95- 
96-5. 执行可执行文件。
97- 
98- ```
99- ./demo
100- ```
101- 
102-## 程序的编译与执行<a name="section313754223513"></a>
103- 
104-通过毕昇编译器可以将算子源文件（以.asc为后缀）编译为当前平台的可执行文件。此外，也支持使用-x asc编译选项编译以.cpp/.c等为后缀的C++/C源文件。
105- 
106-```
107- # bisheng -x asc [算子源文件] -o [输出产物名称] --npu-arch=[NPU架构版本号] --enable-simt，常见参数顺序与g++保持一致。
108-# C++源文件
109-bisheng -x asc add_custom.cpp -o add_custom --npu-arch=dav-xxxx --enable-simt
110-# 后缀为.asc的算子源文件
111-bisheng -x asc add_custom.asc -o add_custom --npu-arch=dav-xxxx --enable-simt
112-```
113- 
114-生成的可执行文件可通过如下方式执行：
115- 
116-```
117-./add_custom
118-```
119- 
120-在命令行编译场景下，可以按需链接需要的库文件，编译时会默认链接[表1](通过CMake编译.md#table201231542115513)中列出的库文件。
121- 
Adocs/guide/编程指南/编译与运行/异步执行.md+66-0
@@ -0,0 +1,66 @@
1+# 异步执行
2+当开发者完成一个Kernel核函数的编写后，需要通过特定的下发/启动语法将其提交到指定的Stream（计算流）上执行。例如：
3+```c++
4+ kernel<<<numBlocks, ubufDynSize, stream>>>(参数列表);
5+ // 或使用对应的API：aclrtLaunchKernel*(...)
6+```
7+这些核Kernel核函数都是通过CANN Runtime完成Kernel的加载与运行。CANN Runtime是CANN软件栈中负责驱动硬件执行与管理AI计算任务的核心组件，它通过提供Device、Memory、Context、Stream、Kernel管理等API，使得上层应用和框架能够高效利用AI处理器的硬件计算资源。
8+ 
9+完成Kernel加载和运行的主要流程分以下几步：
10+```txt
11+ 1. 初始化：aclInit。
12+ 2. 运行时资源申请：通过aclrtSetDevice和aclrtCreateStream分别申请Device、Stream运行管理资源。
13+ 3. 如果需要Device内存参与计算，使用aclrtMalloc分配Device内存，通过aclrtMemcpy或aclrtMemcpyAsync将Host侧内存拷贝到Device侧。
14+ 4. 通过kernel<<<numBlocks, ubufDynSize, stream>>>(参数列表) 调用核函数，完成核函数的异步执行。
15+ 5. 执行核函数后，需要同步等待核函数的执行完成。
16+ 6. 如果Host需要Device内存的计算结果，使用aclrtMemcpy或aclrtMemcpyAsync将Device侧内存拷贝到Host侧。
17+ 7. 资源释放：通过aclrtDestroyStream和aclrtResetDevice分别释放Stream、Device运行管理资源。
18+ 8. 去初始化：aclFinalize。
19+```
20+![](../../figures/NPU侧运行验证步骤.png)
21+ 
22+在这个过程中我们提到了Stream创建与销毁，内存申请与拷贝，异步调用与同步等待。在本节的其余部分，我们将解释这些CANN Runtime的元素和对应的API。
23+ 
24+## Stream管理
25+Stream（计算流）本质上是一个有序的任务队列，为开发者提供了一种表示“先序执行”的抽象机制，程序可向队列中提交各类操作，例如内存拷贝、核函数下发等，所有任务将按入队顺序依次执行。单个计算流内，队列头部任务优先执行，执行完成后自动出队；后续队列任务依次前移，等待调度执行。同一个计算流中的任务具备串行执行特性，严格遵循入队先后顺序运行。
26+ 
27+应用程序可同时创建并使用多个计算流。多流场景下，运行时会根据AI 处理器的硬件资源占用状态，从存有待处理任务的计算流中择优调度任务执行。开发者可为计算流配置优先级，该优先级将作为运行时调度的参考依据，用以调整任务调度策略，但无法强制保证固定的执行顺序。
28+ 
29+绑定在指定计算流内的API接口调用、核函数下发操作，相对于主机CPU线程均为异步执行。
30+应用程序可通过等待目标计算流清空所有任务完成流级同步，也可在整个设备维度完成全局同步。
31+ 
32+CANN Runtime内置默认计算流，所有未显式指定计算流的运算操作、核函数下发任务，都会默认排入该流中执行。绝大多数未手动指定计算流的业务代码，都在隐式使用默认计算流。
33+ 
34+- **创建与销毁Stream**： `aclrtCreateStream(&stream)`与`aclrtDestroyStream(stream)`。
35+- **指定Stream执行核函数**： `kernel<<<numBlocks, 0, stream>>>(...)`。
36+- **指定Stream异步拷贝**： `aclrtMemcpyAsync(..., stream)`。
37+- **指定Stream同步**：`aclrtSynchronizeStream(stream)`。
38+ 
39+## Memory（内存）管理
40+在异构计算架构中，系统由主机（Host）和设备（Device）组成。Host和Device各自拥有独立的内存，Host内存是指AI处理器所在服务器的主机内存（即CPU内存），而Device内存则是指AI处理器自带的设备内存。
41+ 
42+内存管理中要做好的两件事是：
43+1. **可以访问内存**：Runtime提供了一套内存管理API，使开发者能够高效便捷地编写应用程序中的内存管理代码。由于Host和Device的内存相互独立，Runtime提供了专门的接口来分别申请和释放Host内存及Device内存。例如，申请和释放Host内存的接口为aclrtMallocHost和aclrtFreeHost，而申请和释放Device内存的接口为aclrtMalloc和aclrtFree。
44+2. **高效访问内存**：算子在Device上执行过程中，访问Device上的Device内存数据性能更高。为此，Runtime提供了Host与Device之间互相拷贝内存的接口，支持同步和异步方式，例如aclrtMemcpy和aclrtMemcpyAsync等，以便开发者更好地规划数据的存储与访问。
45+ 
46+## 异步调用
47+为了能够更好的发挥CPU和AI处理器的硬件利用率，CANN Runtime遵循“下发即返回”的异步执行策略，使 Host 在提交 H2D/D2H 拷贝或 Kernel 计算任务后立即继续执行后续逻辑，从而让主机侧调度、设备侧计算、主机与设备之间的数据传输尽可能重叠执行，隐藏传输和调度开销，提高端到端吞吐率。
48+ 
49+典型<<<\>\>\>异步调用的场景是：将主机与设备计算操作重叠执行，从而提升主机侧与设备侧计算资源利用率。
50+![ ](../../figures/asynchronous_execution.png "Kernel异步调用")
51+ 
52+ 
53+同样地，将aclrtMemcpy替换成aclrtMemcpyAsync接口，可以实现主机与设备的内存传输与计算操作重叠执行，从而降低甚至消除内存传输带来的性能开销。
54+![ ](../../figures/asynchronous_execution_async.png "内存传输异步并发执行")
55+ 
56+ 
57+ 
58+## 同步等待
59+核函数的调用是异步的，在核函数调用结束之后，控制权立刻返回给主机端，需要在host等待核函数的执行完成，可以通过以下接口强制Host程序等待核函数执行完毕：
60+|同步接口|接口描述|
61+|-|-|
62+|`aclrtSynchronizeDevice()` | 阻塞Host侧当前线程，直到与当前线程绑定的Context所对应的Device中的所有操作都已完成。|
63+|`aclrtSynchronizeStream()` |阻塞Host侧当前线程直到指定Stream中的所有操作都已完成。 |
64+ 
65+## 更多的Runtime高级特性
66+更多的Runtime编程方式及接口，可以参考[《应用开发（C&C++）》](https://hiascend.com/document/redirect/CannCommunityadev)。
Rdocs/guide/编程指南/编译与运行/AI-CPU算子编译.md→docs/guide/编程指南/编译与运行/算子编译/AI-CPU算子编译基本用法.md+2-60
@@ -61,7 +61,7 @@ $bisheng --npu-arch=dav-2201 main.asc -c -o main.asc.o
61$bisheng hello_world.aicpu.o main.asc.o -o demo61$bisheng hello_world.aicpu.o main.asc.o -o demo
62```62```
63 63 
64-上文我们通过一个入门示例介绍了使用bisheng命令行编译生成可执行文件的示例。除此之外，使用bisheng命令行也支持编译生成AI CPU算子的动态库与静态库，用户可在asc代码中通过内核调用符<<<...\>\>\>调用AI CPU算子的核函数，并在编译asc代码源文件生成可执行文件的时候，链接AI CPU动态库或者静态库，注意：若单独编译AI CPU算子代码生成动态库、静态库时，需要手动链接[表2](AI-Core-SIMT编译/通过CMake编译.md#table201231542115513)。64+上文我们通过一个入门示例介绍了使用bisheng命令行编译生成可执行文件的示例。除此之外，使用bisheng命令行也支持编译生成AI CPU算子的动态库与静态库，用户可在asc代码中通过内核调用符<<<...\>\>\>调用AI CPU算子的核函数，并在编译asc代码源文件生成可执行文件的时候，链接AI CPU动态库或者静态库，注意：若单独编译AI CPU算子代码生成动态库、静态库时，需要手动链接[表2](AI-Core算子编译基本用法.md#table201231542115513)。
65 65 
66- 编译生成算子动态库66- 编译生成算子动态库
67 67 
@@ -173,7 +173,7 @@ AI CPU算子常用的编译选项说明如下：
173</tbody>173</tbody>
174</table>174</table>
175 175 
176-## 通过CMake编译<a name="section1121825118533"></a>176+## CMake方式编译<a name="section1121825118533"></a>
177 177 
178项目中可以使用CMake来更简便地使用毕昇编译器编译AI CPU算子，生成可执行文件、动态库、静态库或二进制文件。178项目中可以使用CMake来更简便地使用毕昇编译器编译AI CPU算子，生成可执行文件、动态库、静态库或二进制文件。
179 179 
@@ -358,61 +358,3 @@ target_compile_options(demo PRIVATE
358</tr>358</tr>
359</tbody>359</tbody>
360</table>360</table>
M
Mmunanhw5月9日

下面的链接库 现在不需要链接了么

AI CPU的链接到AI Core中去感觉有点奇怪

likedislike
W
wk0911
5月9日 评论：
W
wk0911
5月9日 评论：
361- 
362-**表 3** 编译AI CPU算子需要手动链接的库
363- 
364-<a name="table346945145419"></a>
365-<table><thead align="left"><tr id="row247245105412"><th class="cellrowborder" valign="top" width="23.98%" id="mcps1.2.3.1.1"><p id="p247194525413"><a name="p247194525413"></a><a name="p247194525413"></a>名称</p>
366-</th>
367-<th class="cellrowborder" valign="top" width="76.02%" id="mcps1.2.3.1.2"><p id="p847154515416"><a name="p847154515416"></a><a name="p847154515416"></a>作用描述</p>
368-</th>
369-</tr>
370-</thead>
371-<tbody><tr id="row1747184585416"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p1047154517544"><a name="p1047154517544"></a><a name="p1047154517544"></a>libascendc_runtime.a</p>
372-</td>
373-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p1647545185419"><a name="p1647545185419"></a><a name="p1647545185419"></a>Ascend C算子参数等组装库。</p>
374-</td>
375-</tr>
376-<tr id="row2471459548"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p194724511542"><a name="p194724511542"></a><a name="p194724511542"></a>libruntime.so</p>
377-</td>
378-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p20471845185417"><a name="p20471845185417"></a><a name="p20471845185417"></a>Runtime运行库。</p>
379-</td>
380-</tr>
381-<tr id="row124734515546"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p147134512548"><a name="p147134512548"></a><a name="p147134512548"></a>libprofapi.so</p>
382-</td>
383-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p14719452541"><a name="p14719452541"></a><a name="p14719452541"></a>Ascend C算子运行性能数据采集库。</p>
384-</td>
385-</tr>
386-<tr id="row1647345105420"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p104784517542"><a name="p104784517542"></a><a name="p104784517542"></a>libunified_dlog.so</p>
387-</td>
388-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p147184545410"><a name="p147184545410"></a><a name="p147184545410"></a>CANN日志收集库。</p>
389-</td>
390-</tr>
391-<tr id="row54717452542"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p1247124511546"><a name="p1247124511546"></a><a name="p1247124511546"></a>libmmpa.so</p>
392-</td>
393-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p34774520544"><a name="p34774520544"></a><a name="p34774520544"></a>CANN系统接口库。</p>
394-</td>
395-</tr>
396-<tr id="row247114520549"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p8471645165412"><a name="p8471645165412"></a><a name="p8471645165412"></a>libascend_dump.so</p>
397-</td>
398-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p1947164513547"><a name="p1947164513547"></a><a name="p1947164513547"></a>CANN维测信息库。</p>
399-</td>
400-</tr>
401-<tr id="row747114515415"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p6473454549"><a name="p6473454549"></a><a name="p6473454549"></a>libc_sec.so</p>
402-</td>
403-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p1847204512547"><a name="p1847204512547"></a><a name="p1847204512547"></a>CANN安全函数库。</p>
404-</td>
405-</tr>
406-<tr id="row3472045145410"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p124819454540"><a name="p124819454540"></a><a name="p124819454540"></a>liberror_manager.so</p>
407-</td>
408-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p2484452543"><a name="p2484452543"></a><a name="p2484452543"></a>CANN错误信息管理库。</p>
409-</td>
410-</tr>
411-<tr id="row7487453546"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p84884585415"><a name="p84884585415"></a><a name="p84884585415"></a>libascendcl.so</p>
412-</td>
413-<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p124817452541"><a name="p124817452541"></a><a name="p124817452541"></a>acl相关接口库。</p>
414-</td>
415-</tr>
416-</tbody>
417-</table>
418- 
Adocs/guide/编程指南/编译与运行/算子编译/AI-Core算子编译基本用法.md+381-0
@@ -0,0 +1,381 @@
1+# AI Core编译基本用法<a name="section229217121411"></a>
2+当开发者完成一个Kernel核函数的编写，并在Host侧通过<<<\>\>\>方式完成核函数的启动调用后，就需要进行算子源码编译与算子运行，接下来我们介绍通过bisheng编译器完成算子编译的基本用法。
3+ 
4+## AI Core SIMD编译
5+使用bisheng编译Ascend C源文件的基本命令如下，其中--npu-arch=dav-\<npu architecture\>用于指定AI处理器架构版本：
6+```shell
7+ bisheng <source_file>.asc -o <output_file> --npu-arch=dav-<npu architecture>
8+```
9+ 
10+bisheng可以通过-I \<path\>指定头文件路径、-L \<path\> 指定链接库路径、-l \<library\>指定需要链接的动态库或静态库、-D\<macro\>=\<value\>指定宏定义，示例如下：
11+```shell
12+ bisheng example.asc -I <path_to_include> -L <path_to_library> -l <library> -o <output_file> --npu-arch=dav-<npu architecture>
13+```
14+ 
15+AI Core SIMD的基本编译流程如下：Host代码使用Host编译器编译成Host二进制，AI Core SIMD代码分为Cube代码和Vector代码，需分别编译成Cube二进制和Vector二进制，先将Cube二进制和Vector二进制链接成Fatbin文件，再与Host二进制合并生成可执行二进制。
16+![ ](../../../figures/aicore_compilation.png "aicore 编译流程示意图")
M
Mmunanhw5月11日

这个图配套的文字说明在哪里

likedislike
17+ 
18+- 异构编译
19+ ```shell
20+ // ----- example.asc -----
21+ // 指定核函数在AI Core上执行
22+ __global__ __aicore__ kernel()
23+ 
24+ bisheng -c example.asc -o example.o --npu-arch=dav-xxxx
25+ bisheng -c main.cpp -o main.o
26+ bisheng example.o main.o -o main
27+ // 或
28+ bisheng main.cpp example.asc -o main --npu-arch=dav-xxxx
29+ ```
30+ 
31+- 单独编译
32+ 
33+ bisheng异构编译默认采用的是全程序编译模式，该模式要求单个源文件X.asc中编译的设备程序没有任何未解析的外部设备函数和变量引用。设备函数可以调用其他编译单元中定义的设备函数或访问其中定义的设备变量，但必须在bisheng命令行中指定`-dc`编译选项，才能启用不同编译单元间设备代码的链接功能。这种支持跨编译单元链接设备代码与符号的能力，被称为单独编译。
34+ 
35+ 单独编译能实现更灵活的代码组织、缩短编译耗时，并生成体积更小的可执行文件；相较于全程序编译，它会增加一定的编译构建复杂度。此外，设备代码链接可能会影响程序性能，这也是该模式未被设为默认模式的原因。链接时优化（LTO）可有效降低单独编译带来的性能损耗。
36+ 
37+ 单独编译的强制要求：
38+ - 在一个编译单元中定义的非常量设备变量，在其他编译单元中引用时，必须使用extern关键字声明；
39+ - 所有常量设备变量，定义和跨单元引用时都必须使用extern关键字；
40+ - 所有Ascend C 源文件（.asc）都必须使用`-dc`选项编译。
41+ 
42+ 下方示例中，definition.asc 定义了变量和函数，example.asc 对其进行引用；两个文件将分别编译，最终链接为完整可执行文件。
43+ ```c++
44+ // ------- definition.asc -------
45+ extern __gm__ int dev_var = 5;
46+ __aicore__ int device_function();
47+
48+ // ------- example.asc -------
49+ extern __aicore__ int device_varibale;
50+ __aicore__ int device_function();
51+ __global__ __aicore__ void kernel(int *var) {
52+ dev_var = 0;
53+ *var = device_function();
54+ }
55+ ```
56+ 
57+ ```shell
58+ bisheng -dc definition.asc -o definition.o --npu-arch=dav-xxxx
59+ bisheng -dc example.asc -o example.o --npu-arch=dav-xxxx
60+ bisheng definition.o example.o -o program
61+ ```
62+ 
63+- 动态库编译
64+ ```shell
65+ bisheng -shared example.asc -o libexample.so --npu-arch=dav-xxxx
66+ ```
67+ 
68+- 静态库编译
69+ ```shell
70+ bisheng -lib example.asc -o libexample.a --npu-arch=dav-xxxx
71+ ```
72+ 
73+## AI Core SIMT编译
74+ 
75+使用bisheng编译器编译Ascend C SIMT源文件时，需要在AI Core SIMD编译的基本命令上增加`--enable-simt`，示例如下：
76+ 
77+```shell
78+ bisheng <source_file>.asc -o <output_file> --npu-arch=dav-<npu architecture> --enable-simt
79+```
80+ 
81+AI Core SIMT的基本编译流程如下：Host代码使用Host编译器编译成Host二进制，AI Core SIMT代码编译成SIMT二进制，先将SIMT二进制链接成Fatbin文件，再与Host二进制合并生成可执行二进制。
82+![ ](../../../figures/aicore_simt_comilation.png "aicore simt编译流程示意图")
83+ 
84+## 基本编译命令汇总
85+|编译方式|AI Core SIMD编译命令| AI Core SIMT编译命令|
86+|-|-|-|
87+|异构编译| bisheng <source_file>.asc -o \<output_file\> --npu-arch=dav-\<npu architecture\> | bisheng <source_file>.asc -o <output_file> --npu-arch=dav-\<npu architecture\> **--enable-simt**|
88+|独立编译|bisheng **-dc** <source_file>.asc -o \<output_file\>**.o** --npu-arch=dav-\<npu architecture\>| bisheng **-dc** <source_file>.asc -o <output_file>**.o** --npu-arch=dav-\<npu architecture\> **--enable-simt**|
89+|动态库编译|bisheng **-shared** <source_file>.asc -o \<output_file\>**.so** --npu-arch=dav-\<npu architecture\>| bisheng **-shared** <source_file>.asc -o <output_file>**.so** --npu-arch=dav-\<npu architecture\> **--enable-simt**|
90+|静态库编译|bisheng **-lib** <source_file>.asc -o \<output_file\>**.a** --npu-arch=dav-\<npu architecture\>| bisheng **-lib** <source_file>.asc -o <output_file>**.a** --npu-arch=dav-\<npu architecture\> **--enable-simt**|
91+ 
92+ 
93+## 常用的编译选项<a name="ZH-CN_TOPIC_0000002462746461"></a>
94+<a name="zh-cn_topic_0000001856506454_table2716123812212"></a>
95+<table><thead align="left"><tr id="zh-cn_topic_0000001856506454_row13875038172111"><th class="cellrowborder" valign="top" width="33.63636363636363%" id="mcps1.1.4.1.1"><p id="zh-cn_topic_0000001856506454_p38753382219"><a name="zh-cn_topic_0000001856506454_p38753382219"></a><a name="zh-cn_topic_0000001856506454_p38753382219"></a><strong id="zh-cn_topic_0000001856506454_b138752389214"><a name="zh-cn_topic_0000001856506454_b138752389214"></a><a name="zh-cn_topic_0000001856506454_b138752389214"></a>选项</strong></p>
96+</th>
97+<th class="cellrowborder" valign="top" width="9.676767676767676%" id="mcps1.1.4.1.2"><p id="zh-cn_topic_0000001856506454_p0875138132119"><a name="zh-cn_topic_0000001856506454_p0875138132119"></a><a name="zh-cn_topic_0000001856506454_p0875138132119"></a><strong id="zh-cn_topic_0000001856506454_b2875113818214"><a name="zh-cn_topic_0000001856506454_b2875113818214"></a><a name="zh-cn_topic_0000001856506454_b2875113818214"></a>是否必需</strong></p>
98+</th>
99+<th class="cellrowborder" valign="top" width="56.686868686868685%" id="mcps1.1.4.1.3"><p id="zh-cn_topic_0000001856506454_p168751138102115"><a name="zh-cn_topic_0000001856506454_p168751138102115"></a><a name="zh-cn_topic_0000001856506454_p168751138102115"></a><strong id="zh-cn_topic_0000001856506454_b198756387215"><a name="zh-cn_topic_0000001856506454_b198756387215"></a><a name="zh-cn_topic_0000001856506454_b198756387215"></a>说明</strong></p>
100+</th>
101+</tr>
102+</thead>
103+<tbody><tr id="zh-cn_topic_0000001856506454_row18875183852112"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="zh-cn_topic_0000001856506454_p1987573810215"><a name="zh-cn_topic_0000001856506454_p1987573810215"></a><a name="zh-cn_topic_0000001856506454_p1987573810215"></a>-help</p>
104+</td>
105+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="zh-cn_topic_0000001856506454_p587515381215"><a name="zh-cn_topic_0000001856506454_p587515381215"></a><a name="zh-cn_topic_0000001856506454_p587515381215"></a>否</p>
106+</td>
107+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="zh-cn_topic_0000001856506454_p138752383219"><a name="zh-cn_topic_0000001856506454_p138752383219"></a><a name="zh-cn_topic_0000001856506454_p138752383219"></a>查看帮助。</p>
108+</td>
109+</tr>
110+<tr id="zh-cn_topic_0000001856506454_row20875103872119"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p1068416586398"><a name="p1068416586398"></a><a name="p1068416586398"></a>--npu-arch</p>
111+</td>
112+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p106841658103918"><a name="p106841658103918"></a><a name="p106841658103918"></a>是</p>
113+</td>
114+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="p1268425803914"><a name="p1268425803914"></a><a name="p1268425803914"></a>编译时指定的<span id="ph17911124171120"><a name="ph17911124171120"></a><a name="ph17911124171120"></a>AI处理器</span>架构，取值为dav-&lt;arch-version&gt;，其中&lt;arch-version&gt;为NPU架构版本号，各产品型号对应的架构版本号请通过<a href="../../语言扩展层/SIMD-BuiltIn关键字.md#table65291052154114">对应关系表</a>进行查询。</p>
115+</td>
116+</tr>
117+<tr id="row15492131482712"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p4492414142719"><a name="p4492414142719"></a><a name="p4492414142719"></a>--npu-soc</p>
118+</td>
119+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p15492151411271"><a name="p15492151411271"></a><a name="p15492151411271"></a>否</p>
120+</td>
121+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="p11492114142717"><a name="p11492114142717"></a><a name="p11492114142717"></a>编译时指定的<span id="ph12176116185918"><a name="ph12176116185918"></a><a name="ph12176116185918"></a>AI处理器</span>型号，npu-soc和npu-arch同时配置时，优先使能npu-arch。</p>
122+<p id="p15587811201611"><a name="p15587811201611"></a><a name="p15587811201611"></a><span id="ph1195412562597"><a name="ph1195412562597"></a><a name="ph1195412562597"></a>AI处理器</span>的型号请通过如下方式获取：</p>
123+<a name="ul1124912113117"></a><a name="ul1124912113117"></a><ul id="ul1124912113117"><li>针对如下产品：在安装<span id="zh-cn_topic_0000001265392790_ph196874123168"><a name="zh-cn_topic_0000001265392790_ph196874123168"></a><a name="zh-cn_topic_0000001265392790_ph196874123168"></a>AI处理器</span>的服务器执行<strong id="zh-cn_topic_0000001265392790_b17687612191618"><a name="zh-cn_topic_0000001265392790_b17687612191618"></a><a name="zh-cn_topic_0000001265392790_b17687612191618"></a>npu-smi info</strong>命令进行查询，获取<strong id="zh-cn_topic_0000001265392790_b10161437131915"><a name="zh-cn_topic_0000001265392790_b10161437131915"></a><a name="zh-cn_topic_0000001265392790_b10161437131915"></a>Name</strong>信息。实际配置值为AscendName，例如<strong id="zh-cn_topic_0000001265392790_b16284944181920"><a name="zh-cn_topic_0000001265392790_b16284944181920"></a><a name="zh-cn_topic_0000001265392790_b16284944181920"></a>Name</strong>取值为<em id="zh-cn_topic_0000001265392790_i1478775919179"><a name="zh-cn_topic_0000001265392790_i1478775919179"></a><a name="zh-cn_topic_0000001265392790_i1478775919179"></a>xxxyy</em>，实际配置值为Ascend<em id="zh-cn_topic_0000001265392790_i1678775901719"><a name="zh-cn_topic_0000001265392790_i1678775901719"></a><a name="zh-cn_topic_0000001265392790_i1678775901719"></a>xxxyy</em>。<p id="zh-cn_topic_0000001265392790_p3529538154519"><a name="zh-cn_topic_0000001265392790_p3529538154519"></a><a name="zh-cn_topic_0000001265392790_p3529538154519"></a><span id="zh-cn_topic_0000001265392790_ph1483216010188"><a name="zh-cn_topic_0000001265392790_ph1483216010188"></a><a name="zh-cn_topic_0000001265392790_ph1483216010188"></a><term id="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term11962195213215"><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term11962195213215"></a><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term11962195213215"></a>Atlas A2 训练系列产品</term>/<term id="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term184716139811"><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term184716139811"></a><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term184716139811"></a>Atlas A2 推理系列产品</term></span></p>
124+<p id="zh-cn_topic_0000001265392790_p1193517020467"><a name="zh-cn_topic_0000001265392790_p1193517020467"></a><a name="zh-cn_topic_0000001265392790_p1193517020467"></a><span id="zh-cn_topic_0000001265392790_ph783112021813"><a name="zh-cn_topic_0000001265392790_ph783112021813"></a><a name="zh-cn_topic_0000001265392790_ph783112021813"></a><term id="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term354143892110"><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term354143892110"></a><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term354143892110"></a>Atlas 200I/500 A2 推理产品</term></span></p>
125+<p id="zh-cn_topic_0000001265392790_p78319051815"><a name="zh-cn_topic_0000001265392790_p78319051815"></a><a name="zh-cn_topic_0000001265392790_p78319051815"></a><span id="zh-cn_topic_0000001265392790_ph1383116081815"><a name="zh-cn_topic_0000001265392790_ph1383116081815"></a><a name="zh-cn_topic_0000001265392790_ph1383116081815"></a><term id="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term4363218112215"><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term4363218112215"></a><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term4363218112215"></a>Atlas 推理系列产品</term></span></p>
126+<p id="zh-cn_topic_0000001265392790_p38313021813"><a name="zh-cn_topic_0000001265392790_p38313021813"></a><a name="zh-cn_topic_0000001265392790_p38313021813"></a><span id="zh-cn_topic_0000001265392790_ph58317041819"><a name="zh-cn_topic_0000001265392790_ph58317041819"></a><a name="zh-cn_topic_0000001265392790_ph58317041819"></a><term id="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term71949488213"><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term71949488213"></a><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term71949488213"></a>Atlas 训练系列产品</term></span></p>
127+</li><li>针对如下产品，在安装<span id="zh-cn_topic_0000001265392790_ph17911124171120"><a name="zh-cn_topic_0000001265392790_ph17911124171120"></a><a name="zh-cn_topic_0000001265392790_ph17911124171120"></a>AI处理器</span>的服务器执行<strong id="zh-cn_topic_0000001265392790_zh-cn_topic_0000001264656721_zh-cn_topic_0000001117597244_b206066255591"><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001264656721_zh-cn_topic_0000001117597244_b206066255591"></a><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001264656721_zh-cn_topic_0000001117597244_b206066255591"></a>npu-smi info -t board -i </strong><em id="zh-cn_topic_0000001265392790_zh-cn_topic_0000001264656721_zh-cn_topic_0000001117597244_i16609202515915"><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001264656721_zh-cn_topic_0000001117597244_i16609202515915"></a><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001264656721_zh-cn_topic_0000001117597244_i16609202515915"></a>id</em><strong id="zh-cn_topic_0000001265392790_zh-cn_topic_0000001264656721_zh-cn_topic_0000001117597244_b14358631175910"><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001264656721_zh-cn_topic_0000001117597244_b14358631175910"></a><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001264656721_zh-cn_topic_0000001117597244_b14358631175910"></a> -c </strong><em id="zh-cn_topic_0000001265392790_zh-cn_topic_0000001264656721_zh-cn_topic_0000001117597244_i16269732165915"><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001264656721_zh-cn_topic_0000001117597244_i16269732165915"></a><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001264656721_zh-cn_topic_0000001117597244_i16269732165915"></a>chip_id</em>命令进行查询，获取<strong id="zh-cn_topic_0000001265392790_b11257114917192"><a name="zh-cn_topic_0000001265392790_b11257114917192"></a><a name="zh-cn_topic_0000001265392790_b11257114917192"></a>Chip Name</strong>和<strong id="zh-cn_topic_0000001265392790_b72671651121916"><a name="zh-cn_topic_0000001265392790_b72671651121916"></a><a name="zh-cn_topic_0000001265392790_b72671651121916"></a>NPU Name</strong>信息，实际配置值为Chip Name_NPU Name。例如<strong id="zh-cn_topic_0000001265392790_b13136111611203"><a name="zh-cn_topic_0000001265392790_b13136111611203"></a><a name="zh-cn_topic_0000001265392790_b13136111611203"></a>Chip Name</strong>取值为Ascend<em id="zh-cn_topic_0000001265392790_i68701996189"><a name="zh-cn_topic_0000001265392790_i68701996189"></a><a name="zh-cn_topic_0000001265392790_i68701996189"></a>xxx</em>，<strong id="zh-cn_topic_0000001265392790_b51347352112"><a name="zh-cn_topic_0000001265392790_b51347352112"></a><a name="zh-cn_topic_0000001265392790_b51347352112"></a>NPU Name</strong>取值为1234，实际配置值为Ascend<em id="zh-cn_topic_0000001265392790_i82901912141813"><a name="zh-cn_topic_0000001265392790_i82901912141813"></a><a name="zh-cn_topic_0000001265392790_i82901912141813"></a>xxx</em><em id="zh-cn_topic_0000001265392790_i154501458102213"><a name="zh-cn_topic_0000001265392790_i154501458102213"></a><a name="zh-cn_topic_0000001265392790_i154501458102213"></a>_</em>1234。其中：<a name="zh-cn_topic_0000001265392790_ul2747601334"></a><a name="zh-cn_topic_0000001265392790_ul2747601334"></a><ul id="zh-cn_topic_0000001265392790_ul2747601334"><li>id：设备id，通过<strong id="zh-cn_topic_0000001265392790_b83171930133314"><a name="zh-cn_topic_0000001265392790_b83171930133314"></a><a name="zh-cn_topic_0000001265392790_b83171930133314"></a>npu-smi info -l</strong>命令查出的NPU ID即为设备id。</li><li>chip_id：芯片id，通过<strong id="zh-cn_topic_0000001265392790_b18888204343317"><a name="zh-cn_topic_0000001265392790_b18888204343317"></a><a name="zh-cn_topic_0000001265392790_b18888204343317"></a>npu-smi info -m</strong>命令查出的Chip ID即为芯片id。</li></ul>
128+<p id="zh-cn_topic_0000001265392790_p1790216395447"><a name="zh-cn_topic_0000001265392790_p1790216395447"></a><a name="zh-cn_topic_0000001265392790_p1790216395447"></a><span id="zh-cn_topic_0000001265392790_ph2272194216543"><a name="zh-cn_topic_0000001265392790_ph2272194216543"></a><a name="zh-cn_topic_0000001265392790_ph2272194216543"></a>Ascend 950PR/Ascend 950DT</span></p>
129+<p id="zh-cn_topic_0000001265392790_p12136131554410"><a name="zh-cn_topic_0000001265392790_p12136131554410"></a><a name="zh-cn_topic_0000001265392790_p12136131554410"></a><span id="zh-cn_topic_0000001265392790_ph13754548217"><a name="zh-cn_topic_0000001265392790_ph13754548217"></a><a name="zh-cn_topic_0000001265392790_ph13754548217"></a><term id="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term1253731311225"><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term1253731311225"></a><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term1253731311225"></a>Atlas A3 训练系列产品</term>/<term id="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term131434243115"><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term131434243115"></a><a name="zh-cn_topic_0000001265392790_zh-cn_topic_0000001312391781_term131434243115"></a>Atlas A3 推理系列产品</term></span></p>
130+</li></ul>
131+</td>
132+</tr>
133+<tr id="row463118312408"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p172818333408"><a name="p172818333408"></a><a name="p172818333408"></a>-x</p>
134+</td>
135+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p112810333405"><a name="p112810333405"></a><a name="p112810333405"></a>否</p>
136+</td>
137+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="p1147317623814"><a name="p1147317623814"></a><a name="p1147317623814"></a>指定编译语言, 如：-x asc, 表示指定为Ascend C编程语言。</p>
138+</td>
139+</tr>
140+<tr id="row95881839113910"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p1752513455396"><a name="p1752513455396"></a><a name="p1752513455396"></a>-o &lt;file&gt;</p>
141+</td>
142+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p1052564512393"><a name="p1052564512393"></a><a name="p1052564512393"></a>否</p>
143+</td>
144+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="p14525174511393"><a name="p14525174511393"></a><a name="p14525174511393"></a>指定输出文件的名称和位置。</p>
145+</td>
146+</tr>
147+<tr id="row14244162782319"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p12244627192312"><a name="p12244627192312"></a><a name="p12244627192312"></a>-c</p>
148+</td>
149+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p14244162712317"><a name="p14244162712317"></a><a name="p14244162712317"></a>否</p>
150+</td>
151+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="p1824452712312"><a name="p1824452712312"></a><a name="p1824452712312"></a>编译生成目标文件。</p>
152+</td>
153+</tr>
154+<tr id="row14244162782319"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p12244627192312"><a name="p12244627192312"></a><a name="p12244627192312"></a>-dc</p>
155+</td>
156+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p14244162712317"><a name="p14244162712317"></a><a name="p14244162712317"></a>否</p>
157+</td>
158+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="p1824452712312"><a name="p1824452712312"></a><a name="p1824452712312"></a>编译生成relocatable目标文件。</p>
159+</td>
160+</tr>
161+<tr id="zh-cn_topic_0000001856506454_row11875938112115"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p3893173113412"><a name="p3893173113412"></a><a name="p3893173113412"></a>-shared，--shared</p>
162+</td>
163+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p98931836342"><a name="p98931836342"></a><a name="p98931836342"></a>否</p>
164+</td>
165+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="zh-cn_topic_0000001856506454_p128758389219"><a name="zh-cn_topic_0000001856506454_p128758389219"></a><a name="zh-cn_topic_0000001856506454_p128758389219"></a>编译生成动态链接库。</p>
166+</td>
167+</tr>
168+<tr id="row149091846161211"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p19268319526"><a name="p19268319526"></a><a name="p19268319526"></a>-lib，--cce-build-static-lib</p>
169+</td>
170+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p690994614128"><a name="p690994614128"></a><a name="p690994614128"></a>否</p>
171+</td>
172+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="p5909184615121"><a name="p5909184615121"></a><a name="p5909184615121"></a>编译生成静态链接库。编译器会将Device侧的代码进行编译链接，生成Device侧二进制文件，随后将该文件作为Host侧编译的输入进行编译，最后链接生成静态链接库。</p>
173+</td>
174+</tr>
175+<tr id="row8304638153012"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p183051238103015"><a name="p183051238103015"></a><a name="p183051238103015"></a>-g</p>
176+</td>
177+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p18305153833013"><a name="p18305153833013"></a><a name="p18305153833013"></a>否</p>
178+</td>
179+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="p230573823018"><a name="p230573823018"></a><a name="p230573823018"></a>编译时增加调试信息。</p>
180+</td>
181+</tr>
182+<tr id="row168573353120"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p1085713312317"><a name="p1085713312317"></a><a name="p1085713312317"></a>--sanitizer</p>
183+</td>
184+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p685743143118"><a name="p685743143118"></a><a name="p685743143118"></a>否</p>
185+</td>
186+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="p1961764455616"><a name="p1961764455616"></a><a name="p1961764455616"></a>编译时增加代码正确性校验信息。使用sanitizer选项时，需要同步添加-g选项，且不能在-O0场景下使用。</p>
187+<p id="p485717333110"><a name="p485717333110"></a><a name="p485717333110"></a>注意，启用该选项后GlobalTensor默认使能L2 Cache，无法通过AscendC::SetL2CacheHint接口设置不使能L2 Cache的模式。</p>
188+</td>
189+</tr>
190+<tr id="row236103412553"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p114041048134211"><a name="p114041048134211"></a><a name="p114041048134211"></a>-fPIC</p>
191+</td>
192+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p34047482426"><a name="p34047482426"></a><a name="p34047482426"></a>否</p>
193+</td>
194+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="p940474884218"><a name="p940474884218"></a><a name="p940474884218"></a>告知编译器产生位置无关代码。</p>
195+</td>
196+</tr>
197+<tr id="row1757181163816"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p1854419204220"><a name="p1854419204220"></a><a name="p1854419204220"></a>-O</p>
198+</td>
199+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p2054319164211"><a name="p2054319164211"></a><a name="p2054319164211"></a>否</p>
200+</td>
201+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="zh-cn_topic_0000001856506454_p1187523802118"><a name="zh-cn_topic_0000001856506454_p1187523802118"></a><a name="zh-cn_topic_0000001856506454_p1187523802118"></a>用于指定编译器的优化级别，当前支持-O3，-O2，-O0。</p>
202+</td>
203+</tr>
204+<tr id="row16786622645"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p1584686172114"><a name="p1584686172114"></a><a name="p1584686172114"></a><span>--run-mode=</span>sim</p>
205+</td>
206+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p158467611219"><a name="p158467611219"></a><a name="p158467611219"></a>否</p>
207+</td>
208+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="p08911647152012"><a name="p08911647152012"></a><a name="p08911647152012"></a>sim模式：链接时用户添加仿真模式对应的实现库，实现代码在仿真模式下运行，可以查看仿真相关日志，方便用户性能调试。</p>
209+</td>
210+</tr>
211+<tr id="row109271176132"><td class="cellrowborder" valign="top" width="33.63636363636363%" headers="mcps1.1.4.1.1 "><p id="p1092701741311"><a name="p1092701741311"></a><a name="p1092701741311"></a>--enable-simt</p>
212+</td>
213+<td class="cellrowborder" valign="top" width="9.676767676767676%" headers="mcps1.1.4.1.2 "><p id="p49274178137"><a name="p49274178137"></a><a name="p49274178137"></a>否</p>
214+</td>
215+<td class="cellrowborder" valign="top" width="56.686868686868685%" headers="mcps1.1.4.1.3 "><p id="p139275178134"><a name="p139275178134"></a><a name="p139275178134"></a>SIMT编程场景，指定SIMT方式编译。</p>
216+</td>
217+</tr>
218+</tbody>
219+</table>
220+ 
221+更多的编译命令和用法可以参见[《毕昇编译器用户指南》](https://www.hiascend.com/document/redirect/CannCommunityBiSheng)。
222+ 
223+## CMake方式编译<a name="ZH-CN_TOPIC_0000002428982142"></a>
224+ 
225+项目中可以使用CMake来更简便地使用毕昇编译器编译Ascend C算子，生成可执行文件、动态库、静态库或二进制文件。
226+ 
227+以下是CMake脚本的示例及其核心步骤说明：
228+ 
229+```cmake
230+# 1、find_package(ASC)是CMake中用于查找和配置Ascend C编译工具链的命令
231+find_package(ASC)
232+ 
233+# 2、指定项目支持的语言包括ASC，ASC表示支持使用毕昇编译器对Ascend C编程语言进行编译
234+project(kernel_samples LANGUAGES ASC)
235+ 
236+# 3、使用CMake接口编译可执行文件、动态库、静态库、二进制文件
237+add_executable(demo
238+ add_custom.asc
239+)
240+#.....
241+target_compile_options(demo PRIVATE
242+ # --npu-arch用于指定NPU的架构版本，dav-后为架构版本号，各产品型号对应的架构版本号。
243+ # <COMPILE_LANGUAGE:ASC>:表明该编译选项仅对语言ASC生效
244+ $<$<COMPILE_LANGUAGE:ASC>: --npu-arch=dav-3510>
245+)
246+```
247+ 
248+编译AI Core SIMT代码时，需要在编译选项中增加`--enable-simt`，示例如下，完整样例请参考：[LINK](https://gitcode.com/cann/asc-devkit/tree/master/examples/03_simt_api/00_introduction/00_quickstart/hello_world_simt)。
249+ 
250+```cmake
251+target_compile_options(demo PRIVATE
252+ # --npu-arch用于指定NPU的架构版本，dav-后为架构版本号，各产品型号对应的架构版本号。
253+ # <COMPILE_LANGUAGE:ASC>:表明该编译选项仅对语言ASC生效
254+ # --enable-simt: 表明启动SIMT 方式编译
255+ $<$<COMPILE_LANGUAGE:ASC>: --npu-arch=dav-3510 --enable-simt>
256+)
257+```
258+ 
259+以下是动态库、静态库编译示例，同时展示如何将源文件切换为用语言ASC编译：
260+ 
261+- 编译.asc文件生成动态库
262+ 
263+ ```cmake
264+ # 3、使用CMake接口编译动态库
265+ add_library(kernel_shared SHARED
266+ cube.asc
267+ vector.asc
268+ mix.asc
269+ )
270+ ```
271+ 
272+- 编译.asc文件生成静态库
273+ 
274+ ```cmake
275+ # 3、使用CMake接口编译静态库
276+ add_library(kernel_lib STATIC
277+ cube.asc
278+ vector.asc
279+ mix.asc
280+ )
281+ ```
282+**表1 常用的CMAKE配置变量说明**
283+| 变量名称| 配置说明 |
284+|--|--|
285+|CMAKE_BUILD_TYPE | 编译模式选项，可配置为：“Release”，Release版本，不包含调试信息，编译最终发布的版本。“Debug”，Debug版本，包含调试信息，便于开发者开发和调试。 |
286+|CMAKE_INSTALL_PREFIX | 用于指定CMake执行install时，安装的路径前缀，执行install后编译产物（ascendc_library中指定的target以及对应的头文件）会安装在该路径下。默认路径为当前目录的out目录下。
287+ 
288+## 其他编译相关说明
289+ 
290+### 内置编译宏开关<a name="section57020345148"></a>
291+ 
292+内置编译宏开关列表如下：
293+ - **ASCENDC\_DUMP** 用于控制Dump开关，默认开关打开，开发者调用printf/DumpTensor/assert后会有信息打印（需要注意直调工程的kernel文件内存在host函数，如果在host函数内调用了printf接口，也会触发kernel内的printf相关初始化动作，进而影响kernel的执行性能\)；设置为0后，表示开关关闭。示例如下：
294+ ```shell
295+ # 关闭所有算子的printf打印功能
296+ bisheng <source_file>.asc -o <output_file> --npu-arch=dav-<npu architecture> -DASCENDC_DUMP=0
297+ ```
298+ - **ASCENDC\_DEBUG** 用于控制Ascend C API的调测开关，默认开关关闭；增加该编译宏后，表示开关打开，此时接口内部的assert校验生效，校验不通过会有assert日志打屏。开启该功能会对算子实际运行的性能带来一定影响，通常在调测阶段使用。示例如下：
299+ ```shell
300+ bisheng <source_file>.asc -o <output_file> --npu-arch=dav-<npu architecture> -DASCENDC_DEBUG
301+ ```
302+ 当前ASCENDC\_DEBUG功能支持的产品型号为：
303+ Atlas 推理系列产品
304+ Atlas A2 训练系列产品/Atlas A2 推理系列产品
305+ 
306+ - **ENABLE\_CV\_COMM\_VIA\_SSBUF** 用于控制是否使用SSBuffer以及UB到L1 Buffer的硬通道，在涉及CV通信（AIC和AIV）或使用数据搬运API时需关注此选项。开启该选项可以提高相关API的性能或拓展使用更多功能。默认开关关闭；设置为true后，表示开关打开。示例如下：
307+ ```shell
308+ bisheng <source_file>.asc -o <output_file> --npu-arch=dav-<npu architecture> -DENABLE_CV_COMM_VIA_SSBUF=true
309+ ```
310+ 仅在Ascend 950PR/Ascend 950DT支持该选项。
311+ - 从其它硬件平台移植到此平台的算子，开关默认关闭以保持兼容性。
312+ - 在该平台新开发的算子，以下场景需要打开：使用矩阵计算Matmul高阶API，且使用SetTensorScaleA等接口，这些接口属于Ascend 950PR/Ascend 950DT新增的功能，其内部实现使用了SSBuffer；使用DataCopy接口从UB拷贝数据到L1 Buffer。
313+ 
314+ - **NO\_OVERLAP\_IN\_MULTI\_REPEAT**
315+ 该编译选项用于在没有地址重叠的情况下移除不必要的内存同步指令，以提升性能。针对Ascend 950PR/Ascend 950DT，使用基础API的高维切分计算API时，默认会插入内存同步指令以确保在地址重叠等复杂场景下的数据正确性，但这些同步指令会带来性能开销。在追求极致性能的场景下，如果您可以确定代码在任何情况下都不会发生内存重叠，可以使用此选项。
316+ 
317+ 
318+### 内置链接库<a name="section57020345148"></a>
319+ 
320+毕昇编译器默认链接的库文件，列表如下：
321+<a name="table201231542115513"></a>
322+<table><thead align="left"><tr id="row171231542205510"><th class="cellrowborder" valign="top" width="23.98%" id="mcps1.2.3.1.1"><p id="p11123114295513"><a name="p11123114295513"></a><a name="p11123114295513"></a>名称</p>
323+</th>
324+<th class="cellrowborder" valign="top" width="76.02%" id="mcps1.2.3.1.2"><p id="p1412374225512"><a name="p1412374225512"></a><a name="p1412374225512"></a>作用描述</p>
325+</th>
326+</tr>
327+</thead>
328+<tbody><tr id="row5123842135514"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p1212364212559"><a name="p1212364212559"></a><a name="p1212364212559"></a>libascendc_runtime.a</p>
329+</td>
330+<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p1112394218551"><a name="p1112394218551"></a><a name="p1112394218551"></a>Ascend C算子参数等组装库。</p>
331+</td>
332+</tr>
333+<tr id="row612324285519"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p01231423552"><a name="p01231423552"></a><a name="p01231423552"></a>libruntime.so</p>
334+</td>
335+<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p8123164255511"><a name="p8123164255511"></a><a name="p8123164255511"></a>Runtime运行库。</p>
336+</td>
337+</tr>
338+<tr id="row1612374285512"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p2012315425551"><a name="p2012315425551"></a><a name="p2012315425551"></a>libprofapi.so</p>
339+</td>
340+<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p12123164265514"><a name="p12123164265514"></a><a name="p12123164265514"></a>Ascend C算子运行性能数据采集库。</p>
341+</td>
342+</tr>
343+<tr id="row10123134212552"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p812374235515"><a name="p812374235515"></a><a name="p812374235515"></a>libunified_dlog.so</p>
344+</td>
345+<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p412314426554"><a name="p412314426554"></a><a name="p412314426554"></a>CANN日志收集库。</p>
346+</td>
347+</tr>
348+<tr id="row1012384210552"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p15123104219559"><a name="p15123104219559"></a><a name="p15123104219559"></a>libmmpa.so</p>
349+</td>
350+<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p13123242135519"><a name="p13123242135519"></a><a name="p13123242135519"></a>CANN系统接口库。</p>
351+</td>
352+</tr>
353+<tr id="row17124154245516"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p612484265518"><a name="p612484265518"></a><a name="p612484265518"></a>libascend_dump.so</p>
354+</td>
355+<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p101241842175512"><a name="p101241842175512"></a><a name="p101241842175512"></a>CANN维测信息库。</p>
356+</td>
357+</tr>
358+<tr id="row6124164213551"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p111246426558"><a name="p111246426558"></a><a name="p111246426558"></a>libc_sec.so</p>
359+</td>
360+<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p31241442185512"><a name="p31241442185512"></a><a name="p31241442185512"></a>CANN安全函数库。</p>
361+</td>
362+</tr>
363+<tr id="row171241342175514"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p6124124218556"><a name="p6124124218556"></a><a name="p6124124218556"></a>liberror_manager.so</p>
364+</td>
365+<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p61248424557"><a name="p61248424557"></a><a name="p61248424557"></a>CANN错误信息管理库。</p>
366+</td>
367+</tr>
368+<tr id="row512404213550"><td class="cellrowborder" valign="top" width="23.98%" headers="mcps1.2.3.1.1 "><p id="p151243425553"><a name="p151243425553"></a><a name="p151243425553"></a>libascendcl.so</p>
369+</td>
370+<td class="cellrowborder" valign="top" width="76.02%" headers="mcps1.2.3.1.2 "><p id="p1012424213555"><a name="p1012424213555"></a><a name="p1012424213555"></a>acl相关接口库。</p>
371+</td>
372+</tr>
373+</tbody>
374+</table>
375+ 
376+### 高阶API常用链接库 <a name="section57020345148"></a>
377+ 
378+| 使用场景 |名称 | 动态库路径 |
379+|--|--|--|
380+| 使用高阶API相关的Tiling接口时需要同时链接。 |libtiling_api.a</br> libregister.so</br> libgraph_base.so | ${ASCEND_HOME_PATH}/lib64 |
381+| 使用PlatformAscendC相关硬件平台信息接口时需要链接。 | libplatform.so | ${ASCEND_HOME_PATH}/lib64 |
Adocs/guide/编程指南/编译与运行/算子编译/RTC运行时编译.md+100-0
@@ -0,0 +1,100 @@
1+# RTC<a name="ZH-CN_TOPIC_0000002462620725"></a>
2+ 
3+RTC是Ascend C运行时编译库，通过[aclrtc](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/context/RTC.md)接口，在程序运行时，将中间代码动态编译成目标机器码，提升程序运行性能。
4+ 
5+运行时编译库提供以下核心接口：
6+- aclrtcCreateProg：根据输入参数（字符串形式表达的Ascend C源代码等）创建aclrtcProg程序实例。
7+- aclrtcCompileProg：编译给定的程序，支持用户自定义编译选项，比如指定NPU架构版本号：--npu-arch=dav-2201。支持的编译选项可以参考[《毕昇编译器用户指南》](https://www.hiascend.com/document/redirect/CannCommunityBiSheng)。
8+- aclrtcGetBinDataSize：获取编译后的Device侧二进制数据的大小。
9+- aclrtcGetBinData：获取编译后的Device侧二进制数据。
10+- aclrtcDestroyProg：在编译和执行过程结束后，销毁给定的程序。
11+ 
12+编译完成后需要调用如下接口完成（仅列出核心接口）Kernel加载与执行。完整流程和详细接口说明请参考[《Runtime运行时API》](https://hiascend.com/document/redirect/CannCommunityRuntimeApi)中的“Kernel加载与执行”章节。
13+1. 通过aclrtBinaryLoadFromData接口解析由aclrtcGetBinData接口获取的算子二进制数据。
14+2. 获取核函数句柄并根据核函数句柄操作其参数列表，相关接口包括aclrtBinaryGetFunction（获取核函数句柄）、aclrtKernelArgsInit（初始化参数列表）、aclrtKernelArgsAppend（追加拷贝用户设置的参数值如xDevice, yDevice, zDevice）等。
15+3. 调用aclrtLaunchKernelWithConfig接口，启动对应算子的计算任务。
16+ 
17+ 
18+如下是一个使用aclrtc接口编译并运行Add自定义算子的完整样例：
19+ 
20+ 
21+```c++
22+#include <vector>
23+#include "acl/acl.h"
24+#include "acl/acl_rt_compile.h" // 使用aclrtc接口需要包含的头文件
25+ 
26+#define ASCENDC_CHECK(expr) do { \
27+ aclError ret = (expr); \
28+ if (ret != ACL_SUCCESS) { \
29+ fprintf(stderr, \
30+ "Ascend Error: %s:%d code=%d %s\n",\
31+ __FILE__, __LINE__, \
32+ ret, aclrtRecentErrMsg(ret)); \
33+ } \
34+} while(0)
35+ 
36+const char *src = R""""(
37+ #include "debug/asc_printf.h"
38+ extern "C" __global__ __aicore__ hello_world()
39+ {
40+ printf("hello world!!!\n")
41+ }
42+)"""";
43+ 
44+int main(int argc, char *argv[])
45+{
46+ // aclrtc流程，src为用户Device侧源码，通过aclrtcCreateProg来创建编译程序
47+ aclrtcProg prog;
48+ ASCENDC_CHECK(aclrtcCreateProg(&prog, src, "hello_world.asc", 0, nullptr, nullptr));
49+ 
50+ // aclrtc流程，传入毕昇编译器的编译选项，调用aclrtcCompileProg进行编译
51+ const char *options[] = {
52+ "--npu-arch=dav-2201",
53+ };
54+ int numOptions = sizeof(options) / sizeof(options[0]);
55+ ASCENDC_CHECK(aclrtcCompileProg(prog, numOptions, options));
56+ 
57+ // aclrtc流程，获取Device侧二进制内容和大小
58+ size_t binDataSizeRet;
59+ ASCENDC_CHECK(aclrtcGetBinDataSize(prog, &binDataSizeRet));
60+ std::vector<char> deviceELF(binDataSizeRet);
61+ ASCENDC_CHECK(aclrtcGetBinData(prog, deviceELF.data()));
62+ 
63+ // --------------------------------------------------- aclrt part --------------------------------------------------
64+ ASCENDC_CHECK(aclrtSetDevice(0));
65+ aclrtStream stream = nullptr;
66+ ASCENDC_CHECK(aclrtCreateStream(&stream));
67+
68+ aclrtBinHandle binHandle = nullptr;
69+ aclrtBinaryLoadOptions loadOption;
70+ loadOption.numOpt = 1;
71+ aclrtBinaryLoadOption option;
72+ option.type = ACL_RT_BINARY_LOAD_OPT_LAZY_MAGIC;
73+ option.value.magic = ACL_RT_BINARY_MAGIC_ELF_VECTOR_CORE; // 设置magic值，表示算子在Vector Core上执行
74+ loadOption.options = &option;
75+ ASCENDC_CHECK(aclrtBinaryLoadFromData(deviceELF.data(), binDataSizeRet, &loadOption, &binHandle));
76+
77+ aclrtFuncHandle funcHandle = nullptr;
78+ const char *funcName = "add_custom";
79+ ASCENDC_CHECK(aclrtBinaryGetFunction(binHandle, funcName, &funcHandle));
80+ 
81+ // 核函数入口
82+ uint32_t numBlocks = 8;
83+ ASCENDC_CHECK(aclrtLaunchKernelWithConfig(funcHandle, numBlocks, stream, nullptr, argsHandle, nullptr));
84+ ASCENDC_CHECK(aclrtSynchronizeStream(stream));
85+ ASCENDC_CHECK(aclrtBinaryUnLoad(binHandle));
86+ ASCENDC_CHECK(aclrtDestroyStream(stream));
87+ ASCENDC_CHECK(aclrtResetDevice(deviceId));
88+ ASCENDC_CHECK(aclFinalize());
89+ 
90+ // 编译和运行均已结束，销毁程序
91+ ASCENDC_CHECK(aclrtcDestroyProg(&prog));
92+ return 0;
93+}
94+```
95+ 
96+编译命令如下，编译时需要设置-I\$\{ASCEND_HOME_PATH\}/include，用于找到aclrtc相关头文件，并设置-L\$\{ASCEND_HOME_PATH\}/lib64链接alc_rtc动态库。
97+ 
98+```
99+g++ add_custom.cpp -I${ASCEND_HOME_PATH}/include -L${ASCEND_HOME_PATH}/lib64 -lascendcl -lacl_rtc -o main
100+```
Adocs/guide/编程指南/编译与运行/算子编译/毕昇编译器.md+20-0
@@ -0,0 +1,20 @@
1+# 毕昇编译器简介<a name="ZH-CN_TOPIC_0000002457558450"></a>
2+ 
3+毕昇编译器是专为昇腾AI处理器设计的异构编译器，原生支持Ascend C算子代码的编译，可将用户编写的算子程序高效转换为二进制可执行文件或动态库。编译器的可执行文件名为bisheng，支持x86、aarch64等多种主机平台，并直接提供对设备侧AI Core架构的指令集编译支持。借助毕昇编译器，开发者可以更高效、便捷地完成面向昇腾平台的算子开发与优化工作。
4+ 
5+目前，编译器支持如下型号：
6+- Ascend 950PR/Ascend 950DT
7+- Atlas A3 训练系列产品/Atlas A3 推理系列产品
8+- Atlas A2 训练系列产品/Atlas A2 推理系列产品
9+- Atlas 推理系列产品
10+ 
11+毕昇编译器支持的源码文件核头文件如下：
12+| 文件扩展名 | 描述 | 内容 |
13+|--|--|--|
14+|.c| C源码文件 | Host-only代码 |
15+|.cpp, .cc, .cxx | C++源码文件| Host-only代码 |
16+|.h, .hpp, .hh, .hxx | C/C++ 头文件| Device代码， Host代码， Host/Device混合代码 |
17+|.asc | Ascend C源码文件| Device代码， Host代码， Host/Device混合代码 |
18+ 
19+- 异构编译场景中的编程相关约束请参考[约束说明](约束说明.md)。
20+- 如果需要了解更多的毕昇编译器的基础知识，详细内容请参考[《毕昇编译器用户指南》](https://www.hiascend.com/document/redirect/CannCommunityBiSheng)。
Rdocs/guide/编程指南/编译与运行/AI-Core-SIMD编译/约束说明.md→docs/guide/编程指南/编译与运行/算子编译/约束说明.md+9-8
@@ -4,25 +4,26 @@
4 - 特别地，针对如下型号，无论是否是同一个编译单元多个核函数的场景，均不支持在开发者未设置Kernel类型时进行自动推导。建议开发者手动设置Kernel类型。4 - 特别地，针对如下型号，无论是否是同一个编译单元多个核函数的场景，均不支持在开发者未设置Kernel类型时进行自动推导。建议开发者手动设置Kernel类型。
5 - Ascend 950PR/Ascend 950DT5 - Ascend 950PR/Ascend 950DT
6 - Atlas 推理系列产品6 - Atlas 推理系列产品
7+</br>
7 8 
8- - 针对Atlas 推理系列产品 ，暂不支持设置Kernel类型为KERNEL\_TYPE\_MIX\_VECTOR\_CORE。9+- Kernel函数推荐使用 \_\_cube\_\_ / \_\_vector\_\_ / \_\_mix\_\_(cube, vec) 进行标记，KERNEL\_TASK\_TYPE\_DEFAULT接口即将废弃。
M
Mmunanhw5月9日

为什么删掉了Atlas 推理系列产品 的限定

likedislike
W
wk0911
5月9日 评论：
10+</br>
9 11 
10-- KERNEL\_TASK\_TYPE\_DEFAULT接口需在核函数中进行调用。
11- 纯Scalar算子无法实现自动推导12- 纯Scalar算子无法实现自动推导
12- 
13 需手动标记Kernel函数类型，推荐设置为纯Vector类型，添加\_\_vector\_\_ attribute进行标记：13 需手动标记Kernel函数类型，推荐设置为纯Vector类型，添加\_\_vector\_\_ attribute进行标记：
14 14 
15 ```15 ```
16 __global__ __vector__ __aicore__ void func0(__gm__ uint8* Addr) {16 __global__ __vector__ __aicore__ void func0(__gm__ uint8* Addr) {
17 Addr[1] = Addr[0];17 Addr[1] = Addr[0];
18 AscendC::printf("Hello world");18 AscendC::printf("Hello world");
19- } 19+ }
20 ```20 ```
21-- bfloat16\_t等数据类型在Host侧不支持，使用这些数据类型时，Host和Device不能写在同一个实现文件里。Host侧不支持的数据类型如下：21+</br>
22 22 
23- Ascend 950PR/Ascend 950DT：bfloat16\_t、hifloat8\_t、fp8\_e5m2\_t、fp8\_e4m3fn\_t、fp8\_e8m0\_t、fp4x2\_e2m1\_t、fp4x2\_e1m2\_t、int4x2\_t。23+- bfloat16\_t等数据类型在 Host 端仅支持以 C++ 模板的形式进行定义与声明，具体数据类型如下：
24 24 
25- Atlas A2 训练系列产品/Atlas A2 推理系列产品：bfloat16\_t。25+ **Ascend 950PR/Ascend 950DT**：bfloat16\_t、hifloat8\_t、fp8\_e5m2\_t、fp8\_e4m3fn\_t、fp8\_e8m0\_t、fp4x2\_e2m1\_t、fp4x2\_e1m2\_t、int4x2\_t。
26 26 
27- Atlas A3 训练系列产品/Atlas A3 推理系列产品：bfloat16\_t。27+ **Atlas A2 训练系列产品/Atlas A2 推理系列产品**：bfloat16\_t。
28 28 
29+ **Atlas A3 训练系列产品/Atlas A3 推理系列产品**：bfloat16\_t。
Ddocs/guide/编程指南/编译与运行/算子运行.md+0-33
@@ -1,33 +0,0 @@
1-# 算子运行<a name="ZH-CN_TOPIC_0000002509833413"></a>
2- 
3-算子的计算算法实现通过Ascend C API来完成，而算子的加载调用则使用Runtime API来完成。本章节将结合核函数调用介绍CANN软件栈中Ascend C算子运行时常用的Runtime接口。Runtime接口更多信息与细节可以参考[《Runtime运行时API》](https://hiascend.com/document/redirect/CannCommunityRuntimeApi)。
4- 
5-## 加载和运行代码<a name="section16306163203216"></a>
6- 
7-加载和运行算子时，需要使用Runtime API，完成运行时管理和配置。主要流程和使用到的API如下：
8- 
9-1. 初始化：aclInit。
10-2. 运行时资源申请：通过aclrtSetDevice和aclrtCreateStream分别申请Device、Stream运行管理资源。
11-3. 使用aclrtMallocHost分配Host内存，并进行数据初始化。
12-4. 使用aclrtMalloc分配Device内存，并通过aclrtMemcpy将数据从Host上拷贝到Device上，参与核函数计算。
13-5. 使用<<<\>\>\>调用算子核函数。
14-6. 执行核函数后，将Device上的运算结果拷贝回Host。
15-7. 异步等待核函数执行完成：aclrtSynchronizeStream。
16-8. 资源释放：通过aclrtDestroyStream和aclrtResetDevice分别释放Stream、Device运行管理资源。
17-9. 去初始化：aclFinalize。
18- 
19-![](../../figures/NPU运行验证.png)
20- 
21-## Kernel加载与执行的更多方式<a name="section45874443418"></a>
22- 
23-Kernel的加载与执行也可以通过二进制加载方式实现，这是最底层的接口实现方式。内核调用符<<<...\>\>\>为对底层接口的封装实现。使用时需要bisheng命令行编译将算子源文件编译为二进制.o文件，再通过aclrtLaunchKernelWithConfig等Kernel加载与执行接口完成算子调用。
24- 
25-- Kernel加载与执行接口的具体说明请参考[《Runtime运行时API》](https://hiascend.com/document/redirect/CannCommunityRuntimeApi)中的"Kernel加载与执行"章节。
26-- 关于更多bisheng命令行编译选项的使用介绍，SIMD编程、SIMD与SIMT混合编程场景请参考[常用的编译选项](AI-Core-SIMD编译/常用的编译选项.md)， SIMT编程场景请参考[SIMT编程常用的编译选项](AI-Core-SIMT编译/常用的编译选项.md)。
27-- SIMD编程、SIMD与SIMT混合编程场景完整样例请参考[Kernel加载与执行（加载二进制）样例](https://gitee.com/ascend/samples/tree/master/operator/ascendc/0_introduction/3_add_kernellaunch/AddKernelInvocationAcl)，SIMT编程场景的样例请参考[LINK](https://gitcode.com/cann/asc-devkit/tree/master/examples/03_simt_api/00_introduction/01_add)。
28- 
29->[!NOTE]说明
30->核函数的调用是异步的，核函数的调用结束后，控制权立刻返回给主机端，可以调用以下**aclrtSynchronizeStream**函数来强制主机端程序等待所有核函数执行完毕。
31->```
32->aclError aclrtSynchronizeStream(aclrtStream stream);
33->```
Mdocs/guide/编程指南/调试调优/功能调试/CPU域孪生调试.md+3-4
@@ -2,7 +2,7 @@
2 2 
3本节介绍CPU域调试的方法：CPU侧验证核函数，gdb调试、使用printf命令打印。当前SIMT编程场景不支持。3本节介绍CPU域调试的方法：CPU侧验证核函数，gdb调试、使用printf命令打印。当前SIMT编程场景不支持。
4 4 
5->[!NOTE]说明 5+>[!NOTE] 说明
6>CPU调测过程中，配置日志相关环境变量，可以记录程序的运行过程及异常信息，有助于开发者进行功能调测。6>CPU调测过程中，配置日志相关环境变量，可以记录程序的运行过程及异常信息，有助于开发者进行功能调测。
7>关于环境变量的使用约束以及详细说明，可参见《环境变量参考》中“辅助功能 \> 日志”章节。7>关于环境变量的使用约束以及详细说明，可参见《环境变量参考》中“辅助功能 \> 日志”章节。
8 8 
@@ -35,7 +35,7 @@
35 35 
36 其他CMakeLists.txt项目配置[通过CMake编译](../../编译与运行/AI-Core-SIMD编译/通过CMake编译.md)进行编写。36 其他CMakeLists.txt项目配置[通过CMake编译](../../编译与运行/AI-Core-SIMD编译/通过CMake编译.md)进行编写。
37 37 
38->[!NOTE]说明 38+>[!NOTE] 说明
39>为了实现CPU域与NPU域代码归一，框架在CPU域中仅对部分acl接口进行适配，开发者在使用CPU域调测功能时，仅支持使用如下acl接口，并且不支持用户自行链接**ascendcl库**：39>为了实现CPU域与NPU域代码归一，框架在CPU域中仅对部分acl接口进行适配，开发者在使用CPU域调测功能时，仅支持使用如下acl接口，并且不支持用户自行链接**ascendcl库**：
40>- 有实际功能接口，支持CPU域调用40>- 有实际功能接口，支持CPU域调用
41> - aclDataTypeSize、aclFloat16ToFloat、aclFloatToFloat16。41> - aclDataTypeSize、aclFloat16ToFloat、aclFloatToFloat16。
@@ -48,7 +48,7 @@
48> - 运行时管理48> - 运行时管理
49> aclrtSetDevice、aclrtResetDevice、aclrtCreateStream、aclrtCreateStreamWithConfig、aclrtDestroyStream、aclrtDestroyStreamForce、aclrtSynchronizeStream、aclrtCreateContext、aclrtDestroyContext。49> aclrtSetDevice、aclrtResetDevice、aclrtCreateStream、aclrtCreateStreamWithConfig、aclrtDestroyStream、aclrtDestroyStreamForce、aclrtSynchronizeStream、aclrtCreateContext、aclrtDestroyContext。
50 50 
51-> [!CAUTION]注意51+> [!CAUTION] 注意
52> 在CPU孪生调试模式下，编译器不识别Ascend C提供的[SIMD-BuiltIn关键字](../../语言扩展层/SIMD-BuiltIn关键字.md)及[SIMT-BuiltIn关键字](../../语言扩展层/SIMT-BuiltIn关键字.md)，包括52> 在CPU孪生调试模式下，编译器不识别Ascend C提供的[SIMD-BuiltIn关键字](../../语言扩展层/SIMD-BuiltIn关键字.md)及[SIMT-BuiltIn关键字](../../语言扩展层/SIMT-BuiltIn关键字.md)，包括
53> - **函数执行空间限定符**：如 `__aicore__`、`__global__`、`__host__` 等;53> - **函数执行空间限定符**：如 `__aicore__`、`__global__`、`__host__` 等;
54> - **函数标记宏**：如 `__simd_vf__`、`__simd_callee__`、`__simt_vf__` 等；54> - **函数标记宏**：如 `__simd_vf__`、`__simd_callee__`、`__simt_vf__` 等；
@@ -171,4 +171,3 @@ quit
171printf("xLocal size: %d\n", xLocal.GetSize()); 171printf("xLocal size: %d\n", xLocal.GetSize());
172printf("tileLength: %d\n", tileLength);172printf("tileLength: %d\n", tileLength);
173```173```
174- 
Mdocs/guide/编程指南/调试调优/功能调试/NPU域上板调试.md+107-12
@@ -12,7 +12,7 @@ printf示例如下，printf接口的使用说明和具体约束请参考[printf]
12printf("fmt string %d", 0x123);12printf("fmt string %d", 0x123);
13```13```
14 14 
15->[!NOTE]说明 15+>[!NOTE]说明
16>printf接口打印功能会对算子实际运行的性能带来一定影响，通常在调测阶段使用。开发者可以按需关闭打印功能。具体方法请参考[printf](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/context/printf-147.md)。16>printf接口打印功能会对算子实际运行的性能带来一定影响，通常在调测阶段使用。开发者可以按需关闭打印功能。具体方法请参考[printf](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/context/printf-147.md)。
17 17 
18## SIMD编程通过DumpTensor打印进行调试<a name="section16452184823718"></a>18## SIMD编程通过DumpTensor打印进行调试<a name="section16452184823718"></a>
@@ -45,8 +45,8 @@ DumpTensor: desc=5, addr=0, data_type=float16, position=UB, dump_size=32
45 45 
46```46```
47 47 
48->[!NOTE]说明 48+> [!NOTE] 说明
49->DumpTensor接口打印功能会对算子实际运行的性能带来一定影响，通常在调测阶段使用。开发者可以按需关闭打印功能。具体方法请参考[DumpTensor](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/context/DumpTensor.md)。49+> DumpTensor接口打印功能会对算子实际运行的性能带来一定影响，通常在调测阶段使用。开发者可以按需关闭打印功能。具体方法请参考[DumpTensor](https://gitcode.com/cann/asc-devkit/blob/master/docs/api/context/DumpTensor.md)。
50 50 
51## 使用msSanitizer工具进行异常检测<a name="section931475414217"></a>51## 使用msSanitizer工具进行异常检测<a name="section931475414217"></a>
52 52 
@@ -59,19 +59,114 @@ msSanitizer工具是基于AI处理器的异常检测工具，包含了单算子
59 59 
60具体使用方法请参考[《算子开发工具用户指南》](https://hiascend.com/document/redirect/CannCommercialToolOpDev)。60具体使用方法请参考[《算子开发工具用户指南》](https://hiascend.com/document/redirect/CannCommercialToolOpDev)。
61 61 
A
Aascendhjh5月7日

这里要从场景出发，当前算子调试调优有哪些典型的场景，基于场景使用哪些工具解决？而不是纯粹介绍工具本身。

likedislike
62-**该功能仅在如下场景支持**：62+**异常检测程序编译**
63+```shell
64+// bisheng编译时 增加sanitizer命令行
65+bisheng add_custom.asc -o add_custom --sanitizer --lineinfo
66+```
67+**异常检测工具的使用**
68+```shell
69+// 同时开启多种异常检测能力
70+mssanitizer --tool=memcheck --tool=racecheck --tool=initcheck --tool=synccheck ./add_custom
71+```
72+**内存异常报告解析**
63 73 
64-- 通过[基于样例工程完成Kernel直调](../../附录/基于样例工程完成Kernel直调.md)方式调用算子。74+内存检测异常报告会输出多种不同类型的异常信息，以下将对非法读写给出了简单的异常信息示例：
65- 75+```shell
66-- 通过[单算子API调用](../../附录/工程化算子开发/单算子API调用.md)方式调用算子。76+====== ERROR: illegal read of size 224 // 异常的基本信息,包含非法读写的类型以及被非法访问的字节数,非法读写包括read(非法读取)和write(非法写入)
67-- 间接调用单算子API\(aclnnxxx\)接口：Pytorch框架单算子直调的场景。77+====== at 0x12c0c0015000 on GM in add_custom_kernel // 异常发生的内存位置信息，包含发生的核函数名、地址空间与内存地址，此处的内存地址指一次内存访问中的首地址
78+====== in block aiv(0) on device 0 // 异常代码对应Vector核的block索引
79+====== code in pc current 0x77c (serialNo:10) // 当前异常发生的pc指针和调用api行为的序列号
80+====== #0 ${ASCEND_HOME_PATH}/asc/impl/basic_api/dav_c220/kernel_operator_data_copy_impl.h :58:9 // 以下为异常发生代码的调用栈，包含文件名、行号和列号
81+====== #1 ${ASCEND_HOME_PATH}/asc/impl/basic_api/kernel_operator_data_copy_intf_impl.h:58:9
82+====== #2 ${ASCEND_HOME_PATH}/asc/include/basic_api/kernel_operator_data_copy_intf.h:443:5
83+====== #3 illegal_read_and_write/add_custom.cpp:18:5
84+```
85+更多详细信息可以参考[《异常检测(msSanitizer)》](https://hiascend.com/document/redirect/mindstudio-op-error-check)。
68 86 
69## 使用msDebug工具进行算子调试<a name="section2072113416285"></a>87## 使用msDebug工具进行算子调试<a name="section2072113416285"></a>
70 88 
71-msDebug是一款面向昇腾设备的算子调试工具，用于调试NPU侧运行的算子程序，为算子开发人员提供调试手段，当前只支持SIMD编程场景的程序调试，暂不支持SIMT编程场景的程序调试。msDebug工具支持调试所有的昇腾算子，包含Ascend C算子（Vector、Cube以及融合算子）程序。具体功能包括断点设置、打印变量和内存、单步调试、中断运行、核切换、检查程序状态、调试信息展示、解析Core dump文件，用户可根据实际情况进行选择。具体使用方法请参考[《算子开发工具用户指南》](https://hiascend.com/document/redirect/CannCommercialToolOpDev)。89+msDebug是一款面向昇腾设备的算子调试工具，用于调试NPU侧运行的算子程序，为算子开发人员提供调试手段，当前只支持SIMD编程场景的程序调试，暂不支持SIMT编程场景的程序调试。msDebug工具支持调试所有的昇腾算子，包含Ascend C算子（Vector、Cube以及融合算子）程序。具体功能包括断点设置、打印变量和内存、单步调试、中断运行、核切换、检查程序状态、调试信息展示、解析Core dump文件，用户可根据实际情况进行选择。
72 90 
73-- 通过[基于样例工程完成Kernel直调](../../附录/基于样例工程完成Kernel直调.md)方式调用算子。91+### msdebug 调试进程
74 92 
75-- 通过[单算子API调用](../../附录/工程化算子开发/单算子API调用.md)方式调用算子。93+#### 调试程序编译
76-- 间接调用单算子API\(aclnnxxx\)接口：Pytorch框架单算子直调的场景。94+```shell
95+// bisheng编译时 增加sanitizer命令行
96+bisheng add_custom.asc -o add_custom -O0 -g
97+```
98+#### 调试程序运行
99+```shell
100+// 使用msdebug进行调试
101+$ msdebug ./add_custom
102+msdebug(MindStudio Debugger) is part of MindStudio Operator-dev Tools.
103+The tool provides developers with a mechanism for debugging Ascend kernels running on actual hardware.
104+This enables developers to debug Ascend kernels without being affected by potential changes brought by simulation and emulation environments.
105+(msdebug) target create "python3"
106+Current executable set to '${INSTALL_DIR}/projects/application' (aarch64).
77 107 
108+(msdebug) b matmul_leakyrelu_kernel.cpp:114
109+Breakpoint 1: where = device_debugdata`_ZN17MatmulLeakyKernelIDhDhffE7CopyOutEj_mix_aiv + 240 at matmul_leakyrelu_kernel.cpp:114:14, address = 0x000000000000ff88
110+(msdebug) run
111+```
112+ 
113+#### 断点设置
114+```shell
115+(msdebug) b matmul_leakyrelu_kernel.cpp:114
116+Breakpoint 1: where = device_debugdata`_ZN17MatmulLeakyKernelIDhDhffE7CopyOutEj_mix_aiv + 240 at matmul_leakyrelu_kernel.cpp:114:14, address = 0x000000000000ff88
117+```
118+ 
119+#### 内存与变量打印
120+```shell
121+# 打印LocalTensor
122+(msdebug) p reluOutLocal
123+(AscendC::LocalTensor<float>) $2 = {
124+ AscendC::BaseLocalTensor<float> = {
125+ address_ = (dataLen = 131072, bufferAddr = 0, bufferHandle = "", logicPos = '\n')
126+ }
127+ shapeInfo_ = {
128+ shapeDim = '\0'
129+ originalShapeDim = '\0'
130+ shape = ([0] = 0, [1] = 1092616192, [2] = 4800, [3] = 1473680, [4] = 0, [5] = 1473888, [6] = 0, [7] = 1471968)
131+ originalShape = ([0] = 0, [1] = 3222199212, [2] = 4800, [3] = 1, [4] = 0, [5] = 1473376, [6] = 0, [7] = 1473376)
132+ dataFormat = ND
133+ }
134+}
135+```
136+ 
137+#### 单步调试
138+```shell
139+(msdebug) s
140+Process 177943 stopped
141+[Switching to focus on Kernel matmul_leakyrelu_custom, CoreId 44, Type aiv]
142+* thread #1, name = 'matmul_leakyrelu', stop reason = step over // 通过回显可查看pc的位置，表示单步成功
143+ frame #0: 0x000000000000f048 device_debugdata`_ZN17MatmulLeakyKernelIDhDhffE10CalcOffsetEiiRK11TCubeTilingRiS4_S4_S4__mix_aiv(this=0x0000000000217b60, blockIdx=0, usedCoreNum=2, tiling=0x0000000000217e28, offsetA=0x00000000002175c8, offsetB=0x00000000002175c4, offsetC=0x00000000002175c0, offsetBias=0x00000000002175bc) at matmul_leakyrelu_kernel.cpp:130:18
144+ 127 offsetA = mCoreIndx * tiling.Ka * tiling.singleCoreM;
145+ 128 offsetB = nCoreIndx * tiling.singleCoreN;
146+ 129 offsetC = mCoreIndx * tiling.N * tiling.singleCoreM + nCoreIndx * tiling.singleCoreN;
147+-> 130 offsetBias = nCoreIndx * tiling.singleCoreN;
148+ 131 }
149+ ```
150+ 
151+#### 调试信息展示
152+ ```shell
153+ (msdebug) ascend info cores
154+ CoreId Type Device Stream Task Block PC stop reason
155+ 12 aic 1 3 0 0 0x12c0c00f03b0 breakpoint 1.2
156+* 44 aiv 1 3 0 0 0x12c0c00f8048 step over //* 代表当前正在运行的核
157+ 45 aiv 1 3 0 0 0x12c0c00f801c breakpoint 1.2
158+ ```
159+ 
160+### 解析异常算子dump文件
161+通过环境变量开启算子异常时Dump出Core文件。
162+```shell
163+# aic_err_detail_dump: 导出AI Core的内存存储、寄存器和调用栈信息。
164+export ASCEND_DUMP_SCENE=aic_err_detail_dump
165+# 指定异常算子Dump信息的存储路径，可配置为绝对路径或执行程序的相对路径。
166+export ASCEND_DUMP_PATH=./
167+```
168+使用msdebug解析core文件：
169+```shell
170+msdebug --core output2/extra-info/data-dump/0/xxx.core add_custom
171+```
172+更多详细信息可以参考[《算子调试（msDebug）》](https://hiascend.com/document/redirect/mindstudio-op-msdebug)。
Mdocs/guide/编程指南/调试调优/性能调优.md+104-52
@@ -2,15 +2,18 @@
2 2 
3## 性能采集与分析工具<a name="section20894115120010"></a>3## 性能采集与分析工具<a name="section20894115120010"></a>
A
Aascendhjh5月7日

1、这里要从场景出发，当前算子调试调优有哪些典型的场景，基于场景使用哪些工具解决？而不是纯粹介绍工具本身。 2、同时要增加链接，导向“算子实践参考”中更深入的算子优化方案和手段。

likedislike
4 4 
5-通过毕昇编译器编译生成可执行程序后，使用msProf工具运行NPU模式下生成的可执行文件，可以采集Ascend C算子在AI处理器上执行的性能数据，进行性能精细调优。5+通过毕昇编译器编译生成可执行程序后，使用msProf工具采集Ascend C算子在AI处理器上执行的性能数据，进行性能精细调优。
6 6 
7-- Profiling性能数据采集：使用msprof工具采集Ascend C算子在AI 处理器上执行的性能数据。7+[msProf工具](https://hiascend.com/document/redirect/CannCommercialToolOpDev)包含msprof op和msprof op simulator两种使用方式，协助用户定位算子内存、算子代码以及算子指令的异常，实现全方位的算子调优。两种使用方式的详细说明请参考表1。
8-- Roofline瓶颈分析：通过msprof op生成的visualize\_data.bin文件可通过MindStudio Insight进行可视化呈现，Roofline瓶颈分析图可构建出处理器的性能模型，然后利用该性能模型快速评估出算子的理论性能极限，协助开发者快速识别瓶颈类型。
9-- 指令流水图分析：通过msprof op simulator生成visualize\_data.bin文件或trace.json文件，并进行可视化呈现。指令流水图以指令维度展示时序关系，并关联调用栈快速定位瓶颈位置。
10 8 
11-性能调优工具的具体使用方法请参考[《算子开发工具用户指南》](https://hiascend.com/document/redirect/CannCommercialToolOpDev)。9+**表1 msprof op和msprof op simulator功能说明表**
12 10 
13-## NPU域上板性能调优<a name="section107378341021"></a>11+| 功能名称 | 适用场景 | 使用方式 | 展示的图形 |
12+|--|--|--| -- |
13+| msprof op | 适用于上板性能分析（实际运行环境中的性能分析），可协助用户定位算子内存和性能瓶颈。 | 直接分析运行中的算子，无需额外配置，适合在板环境中快速定位算子性能问题。 | 计算内存热力图<br>Roofline瓶颈分析图<br>Cache热力图<br>通算流水图<br>算子代码热点图|
14+|msprof op simulator| 适用于仿真性能分析，可协助用户分析算子指令和代码热点问题。|需要参考msprof op simulator配置，配置环境变量（如LD_LIBRARY_PATH）和编译选项（如添加-g生成调试信息），适合在仿真环境中详细分析算子行为。|指令流水图<br>算子代码热点图<br>内存通路吞吐率波形图|
15+ 
16+## NPU域上板性能分析<a name="section107378341021"></a>
14 17 
15算子程序通过毕昇编译器编译生成可执行程序后，可以通过msprof op在NPU上完成性能采集。以SIMD编程场景为例，使用msprof工具采集上板性能数据大致步骤如下：18算子程序通过毕昇编译器编译生成可执行程序后，可以通过msprof op在NPU上完成性能采集。以SIMD编程场景为例，使用msprof工具采集上板性能数据大致步骤如下：
16 19 
@@ -32,6 +35,7 @@
32 35 
33 ```36 ```
34 性能数据文件夹结构示例：37 性能数据文件夹结构示例：
38+ OPPROF_{timestamp}_XXX
35 ├──dump # 原始的性能数据，用户无需关注39 ├──dump # 原始的性能数据，用户无需关注
36 ├──ArithmeticUtilization.csv # cube/vector指令cycle占比，建议优化算子逻辑，减少冗余计算指令40 ├──ArithmeticUtilization.csv # cube/vector指令cycle占比，建议优化算子逻辑，减少冗余计算指令
37 ├──L2Cache.csv # L2 Cache命中率，影响MTE2，建议合理规划数据搬运逻辑，增加命中率41 ├──L2Cache.csv # L2 Cache命中率，影响MTE2，建议合理规划数据搬运逻辑，增加命中率
@@ -43,19 +47,91 @@
43 ├──ResourceConflictRatio.csv # UB上的bank group、bank conflict和资源冲突率在所有指令中的占比， 建议减少/避免对于同一个bank读写冲突或bank group的读读冲突47 ├──ResourceConflictRatio.csv # UB上的bank group、bank conflict和资源冲突率在所有指令中的占比， 建议减少/避免对于同一个bank读写冲突或bank group的读读冲突
44 └──visualize_data.bin # MindStudio Insight呈现文件48 └──visualize_data.bin # MindStudio Insight呈现文件
45 ```49 ```
50+
51+**表2 msprof op文件介绍**
52+<a name="table244174535419"></a>
53+<table><thead align="left"><tr id="row17451145145411"><th class="cellrowborder" valign="top" width="24.54%" id="mcps1.2.3.1.1"><p id="p18451345115414"><a name="p18451345115414"></a><a name="p18451345115414"></a>名称</p>
54+</th>
55+<th class="cellrowborder" valign="top" width="75.46000000000001%" id="mcps1.2.3.1.2"><p id="p4451245115415"><a name="p4451245115415"></a><a name="p4451245115415"></a>说明</p>
56+</th>
57+</tr>
58+</thead>
59+<tbody><tr id="row194514510541"><td class="cellrowborder" valign="top" width="24.54%" headers="mcps1.2.3.1.1 "><p id="p04514519549"><a name="p04514519549"></a><a name="p04514519549"></a>dump文件夹</p>
60+</td>
61+<td class="cellrowborder" valign="top" width="75.46000000000001%" headers="mcps1.2.3.1.2 "><p id="p6451456545"><a name="p6451456545"></a><a name="p6451456545"></a>原始的性能数据，用户无需关注。
62+</td>
63+</tr>
64+<tr id="row74518458542"><td class="cellrowborder" valign="top" width="24.54%" headers="mcps1.2.3.1.1 "><p id="p14520459540"><a name="p14520459540"></a><a name="p14520459540"></a>ArithmeticUtilization.csv</p>
65+</td>
66+<td class="cellrowborder" valign="top" width="75.46000000000001%" headers="mcps1.2.3.1.2 "><p id="p145114515419"><a name="p145114515419"></a><a name="p145114515419"></a>Cube和Vector类型的指令耗时和占比，可参考ArithmeticUtilization（Cube及Vector类型指令耗时和占比）。</p>
67+</td>
68+</tr>
69+<tr id="row124534535416"><td class="cellrowborder" valign="top" width="24.54%" headers="mcps1.2.3.1.1 "><p id="p144516455548"><a name="p144516455548"></a><a name="p144516455548"></a>L2Cache.csv</p>
70+</td>
71+<td class="cellrowborder" valign="top" width="75.46000000000001%" headers="mcps1.2.3.1.2 "><p id="p845134505416"><a name="p845134505416"></a><a name="p845134505416"></a>L2 Cache命中率，可参考L2Cache（L2 Cache命中率）。
72+</td>
73+</tr>
74+<tr id="row124534535416"><td class="cellrowborder" valign="top" width="24.54%" headers="mcps1.2.3.1.1 "><p id="p144516455548"><a name="p144516455548"></a><a name="p144516455548"></a>Memory.csv</p>
75+</td>
76+<td class="cellrowborder" valign="top" width="75.46000000000001%" headers="mcps1.2.3.1.2 "><p id="p845134505416"><a name="p845134505416"></a><a name="p845134505416"></a>UB/L1/L2/主存储器采集内存读写带宽速率，可参考Memory（内存读写带宽速率）。
77+</td>
78+</tr>
46 79 
47-对于SIMT编程场景，只需遵循[AI Core SIMT编译](../编译与运行/AI-Core-SIMT编译/算子编译简介.md)指导进行算子编译，生成可执行文件后，按照上述步骤2和步骤3使用msprof工具执行程序，以获取算子执行的性能数据。80+<tr id="row124534535416"><td class="cellrowborder" valign="top" width="24.54%" headers="mcps1.2.3.1.1 "><p id="p144516455548"><a name="p144516455548"></a><a name="p144516455548"></a>MemoryL0.csv</p>
81+</td>
82+<td class="cellrowborder" valign="top" width="75.46000000000001%" headers="mcps1.2.3.1.2 "><p id="p845134505416"><a name="p845134505416"></a><a name="p845134505416"></a>L0A/L0B/L0C采集内存读写带宽速率，可参考MemoryL0（L0读写带宽速率）。
83+</td>
84+</tr>
48 85 
49-## NPU域性能仿真<a name="section75259502193"></a>86+<tr id="row124534535416"><td class="cellrowborder" valign="top" width="24.54%" headers="mcps1.2.3.1.1 "><p id="p144516455548"><a name="p144516455548"></a><a name="p144516455548"></a>MemoryUB.csv</p>
87+</td>
88+<td class="cellrowborder" valign="top" width="75.46000000000001%" headers="mcps1.2.3.1.2 "><p id="p845134505416"><a name="p845134505416"></a><a name="p845134505416"></a>mte/vector/scalar采集ub读写带宽速率，可参考MemoryUB（UB读写带宽速率）。
89+</td>
90+</tr>
91+ 
92+<tr id="row124534535416"><td class="cellrowborder" valign="top" width="24.54%" headers="mcps1.2.3.1.1 "><p id="p144516455548"><a name="p144516455548"></a><a name="p144516455548"></a>PipeUtilization.csv</p>
93+</td>
94+<td class="cellrowborder" valign="top" width="75.46000000000001%" headers="mcps1.2.3.1.2 "><p id="p845134505416"><a name="p845134505416"></a><a name="p845134505416"></a>采集计算单元和搬运单元耗时和占比，可参考PipeUtilization（计算单元和搬运单元耗时占比）。
95+</td>
96+</tr>
97+ 
98+<tr id="row124534535416"><td class="cellrowborder" valign="top" width="24.54%" headers="mcps1.2.3.1.1 "><p id="p144516455548"><a name="p144516455548"></a><a name="p144516455548"></a>ResourceConflictRatio.csv</p>
99+</td>
100+<td class="cellrowborder" valign="top" width="75.46000000000001%" headers="mcps1.2.3.1.2 "><p id="p845134505416"><a name="p845134505416"></a><a name="p845134505416"></a>UB上的bank group、bank conflict和资源冲突在所有指令中的占比，可参考ResourceConflictRatio（资源冲突占比）。
101+</td>
102+</tr>
103+ 
104+<tr id="row124534535416"><td class="cellrowborder" valign="top" width="24.54%" headers="mcps1.2.3.1.1 "><p id="p144516455548"><a name="p144516455548"></a><a name="p144516455548"></a>OpBasicInfo.csv</p>
105+</td>
106+<td class="cellrowborder" valign="top" width="75.46000000000001%" headers="mcps1.2.3.1.2 "><p id="p845134505416"><a name="p845134505416"></a><a name="p845134505416"></a>算子基础信息，包含算子名称、block dim和耗时等信息，可参考OpBasicInfo（算子基础信息）。
107+</td>
108+</tr>
109+ 
110+<tr id="row124534535416"><td class="cellrowborder" valign="top" width="24.54%" headers="mcps1.2.3.1.1 "><p id="p144516455548"><a name="p144516455548"></a><a name="p144516455548"></a>visualize_data.bin</p>
111+</td>
112+<td class="cellrowborder" valign="top" width="75.46000000000001%" headers="mcps1.2.3.1.2 "><p id="p845134505416"><a name="p845134505416"></a><a name="p845134505416"></a>算子基础信息、计算单元负载、热点函数和Roofline瓶颈分析等信息的可视化呈现文件。visualize_data.bin可通过MindStudio Insight工具进行可视化展示。
113+</td>
114+</tr>
115+<tr id="row124534535416"><td class="cellrowborder" valign="top" width="24.54%" headers="mcps1.2.3.1.1 "><p id="p144516455548"><a name="p144516455548"></a><a name="p144516455548"></a>trace.json</p>
116+</td>
117+<td class="cellrowborder" valign="top" width="75.46000000000001%" headers="mcps1.2.3.1.2 "><p id="p845134505416"><a name="p845134505416"></a><a name="p845134505416"></a>通算流水可视化呈现文件，Chrome浏览器具体请参考通算流水图。
118+</td>
119+</tr>
120+</tbody>
121+</table>
122+ 
123+对于SIMT编程场景，只需遵循[AI Core SIMT编译](../编译与运行/算子编译/AI-Core算子编译基本用法.md)指导进行算子编译，生成可执行文件后，按照上述步骤2和步骤3使用msprof工具执行程序，以获取算子执行的性能数据。
124+ 
125+## NPU域仿真性能分析<a name="section75259502193"></a>
50 126 
51在非昇腾设备上，通过毕昇编译器仿真编译后生成可执行程序，可以通过msprof op simulator完成性能流水仿真。当前仅支持SIMD编程场景，SIMT编程场景不支持。127在非昇腾设备上，通过毕昇编译器仿真编译后生成可执行程序，可以通过msprof op simulator完成性能流水仿真。当前仅支持SIMD编程场景，SIMT编程场景不支持。
52 128 
53- **通过CMake方式仿真编译算子**129- **通过CMake方式仿真编译算子**
54- - 方法一：灵活控制不同的target是否开启仿真编译。修改CMakeList，使用target\_link\_libraries与target\_link\_directories手动配置链接库与路径：130+ - 灵活控制不同的target是否开启仿真编译。修改CMakeList，使用target\_link\_libraries与target\_link\_directories手动配置链接库与路径：
55 131 
56 ```132 ```
57 find_package(ASC REQUIRED)133 find_package(ASC REQUIRED)
58- project(kernel_samples LANGUAGES ASC CXX)134+ project(kernel_samples LANGUAGES ASC)
59 135
60 add_executable(demo136 add_executable(demo
61 add_custom.asc137 add_custom.asc
@@ -82,36 +158,9 @@
82 158 
83 其中仿真库所在目录与NPU架构版本号之间的关系如下，目录名使用**下划线**连接“dav”和架构版本号。159 其中仿真库所在目录与NPU架构版本号之间的关系如下，目录名使用**下划线**连接“dav”和架构版本号。
84 160 
85- **表 1** simulator目录名称与npu-arch关系
86- 
87- <a name="table1871305213304"></a>
88- <table><thead align="left"><tr id="row17714165233011"><th class="cellrowborder" valign="top" width="46%" id="mcps1.2.3.1.1"><p id="p127146529308"><a name="p127146529308"></a><a name="p127146529308"></a>npu-arch</p>
89- </th>
90- <th class="cellrowborder" valign="top" width="54%" id="mcps1.2.3.1.2"><p id="p18714952173010"><a name="p18714952173010"></a><a name="p18714952173010"></a>simulator目录名称</p>
91- </th>
92- </tr>
93- </thead>
94- <tbody><tr id="row1231154754913"><td class="cellrowborder" valign="top" width="46%" headers="mcps1.2.3.1.1 "><p id="p82311747134918"><a name="p82311747134918"></a><a name="p82311747134918"></a>dav-2002</p>
95- </td>
96- <td class="cellrowborder" valign="top" width="54%" headers="mcps1.2.3.1.2 "><p id="p5232164774917"><a name="p5232164774917"></a><a name="p5232164774917"></a>dav_2002</p>
97- </td>
98- </tr>
99- <tr id="row171495215306"><td class="cellrowborder" valign="top" width="46%" headers="mcps1.2.3.1.1 "><p id="p16714052163019"><a name="p16714052163019"></a><a name="p16714052163019"></a>dav-2201</p>
100- </td>
101- <td class="cellrowborder" valign="top" width="54%" headers="mcps1.2.3.1.2 "><p id="p9714155273011"><a name="p9714155273011"></a><a name="p9714155273011"></a>dav_2201</p>
102- </td>
103- </tr>
104- <tr id="row187141852123012"><td class="cellrowborder" valign="top" width="46%" headers="mcps1.2.3.1.1 "><p id="p6714185210308"><a name="p6714185210308"></a><a name="p6714185210308"></a>dav-3510</p>
105- </td>
106- <td class="cellrowborder" valign="top" width="54%" headers="mcps1.2.3.1.2 "><p id="p14714452173013"><a name="p14714452173013"></a><a name="p14714452173013"></a>dav_3510</p>
107- </td>
108- </tr>
109- </tbody>
110- </table>
111- 
112 其中仿真编译所依赖的库介绍如下，开启仿真编译时，需要优先链接，确保优先使用仿真库的符号，防止出现运行时coredump等异常情况。161 其中仿真编译所依赖的库介绍如下，开启仿真编译时，需要优先链接，确保优先使用仿真库的符号，防止出现运行时coredump等异常情况。
113 162 
114- **表 2** 仿真编译依赖库介绍163+ **表3 仿真编译依赖库介绍**
115 164 
116 <a name="table84131311195116"></a>165 <a name="table84131311195116"></a>
117 <table><thead align="left"><tr id="row4413911115113"><th class="cellrowborder" valign="top" width="45.97%" id="mcps1.2.3.1.1"><p id="p1841381120511"><a name="p1841381120511"></a><a name="p1841381120511"></a>名称</p>166 <table><thead align="left"><tr id="row4413911115113"><th class="cellrowborder" valign="top" width="45.97%" id="mcps1.2.3.1.1"><p id="p1841381120511"><a name="p1841381120511"></a><a name="p1841381120511"></a>名称</p>
@@ -133,43 +182,37 @@
133 </tbody>182 </tbody>
134 </table>183 </table>
135 184 
136- - 方法二：通过向CMake传入变量CMAKE\_ASC\_RUN\_MODE和CMAKE\_ASC\_ARCHITECTURES来统一开启仿真编译。命令示例如下：185+ - 通过向CMake传入变量CMAKE\_ASC\_RUN\_MODE和CMAKE\_ASC\_ARCHITECTURES来统一开启仿真编译。命令示例如下：
137 186 
138 sim表示开启仿真编译，dav-后为NPU架构版本号，请根据实际情况进行填写。187 sim表示开启仿真编译，dav-后为NPU架构版本号，请根据实际情况进行填写。
139 188 
140 ```189 ```
141 cmake -B build -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201190 cmake -B build -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201
142- cmake --build build -j
143 ```191 ```
144 192 
145- >[!NOTE]说明 193+ >[!NOTE] 说明
146 >使用命令行往CMake传入变量的方式全局生效，会对CMakeList中所有的target开启sim模式。194 >使用命令行往CMake传入变量的方式全局生效，会对CMakeList中所有的target开启sim模式。
147 195 
148- **通过命令行方式仿真编译算子可执行程序**196- **通过命令行方式仿真编译算子可执行程序**
149 197 
150 ```198 ```
151- # ${INSTALL_DIR}请替换为CANN软件安装后文件存储路径。以root用户安装为例，安装后文件默认存储路径为：/usr/local/Ascend/cann。199+ # 编译算子: bisheng [算子源文件] -o [输出产物名称] --npu-arch=[NPU架构版本号]，--run-mode=sim
152- # 设置simulator目录名称200+ bisheng add_custom.asc -o add_custom --npu-arch=dav-2201 --run-mode=sim
153- export SIMULATOR_FOLDER=dav_2201
154- # 编译算子object: bisheng -c [算子源文件] -o [输出产物名称] --npu-arch=[NPU架构版本号]，--run-mode=sim
155- bisheng -c add_custom.asc -o add_custom.asc.o --npu-arch=dav-2201 --run-mode=sim
156- # 将add算子object转为可执行程序
157- bisheng -Wl,--disable-new-dtags -L${INSTALL_DIR}/tools/simulator/${SIMULATOR_FOLDER}/lib -Wl,-rpath,${INSTALL_DIR}/tools/simulator/${SIMULATOR_FOLDER}/lib -lruntime_camodel -lnpu_drv_camodel -lm -lstdc++ -lascendcl -lascendc_runtime -lprofapi -lunified_dlog -lmmpa -lascend_dump -lc_sec -lerror_manager -lnpu_drv add_custom.asc.o -o add_custom
158- ```
159 201 
160- 编译时除了需要连接仿真库libruntime\_camodel.so和libnpu\_drv.so，还需要链接libascendc\_runtime.a、libruntime.so、libprofapi.so、libunified\_dlog.so、libmmpa.so、libascend\_dump.so、libc\_sec.so、liberror\_manager.so、libascendcl.so（具体说明参考[表2](../编译与运行/AI-Core-SIMD编译/通过CMake编译.md#table201231542115513)）和第三方库libstdc++.so、libm.so。202+ ```
161 203 
162- **性能流水仿真**204- **性能流水仿真**
163 205 
164 使用msprof op simulator并获取仿真数据。206 使用msprof op simulator并获取仿真数据。
165 207 
166 ```208 ```
167- msprof op simulator ./add_custom209+ msprof op simulator ./add_custom --soc-version=Ascendxxxyy
168 ```210 ```
169- 211+
170 仿真数据说明212 仿真数据说明
171 213 
172 ```214 ```
215+ OPPROF_{timestamp}_XXX
173 ├──dump # 原始的性能数据，用户无需关注216 ├──dump # 原始的性能数据，用户无需关注
174 └──simulator # 算子基础信息217 └──simulator # 算子基础信息
175 ├──core0.cubecore0218 ├──core0.cubecore0
@@ -178,4 +221,13 @@
178 ├──trace.json # Edge/Chrome Trace Viewer/Perfetto呈现文件221 ├──trace.json # Edge/Chrome Trace Viewer/Perfetto呈现文件
179 └──visualize_data.bin # MindStudio Insight呈现文件222 └──visualize_data.bin # MindStudio Insight呈现文件
180 ```223 ```
224+**表4 msprof op simulator文件介绍**
225+| 名称 | 说明 |
226+|--|--|
227+| dump文件夹 | 原始仿真生成的dump数据存放文件夹。 |
228+| core*_code_exe.csv | 代码行耗时，*代表0~n核，以便用户快速确定编写的代码中最耗时的部分，可参考代码行耗时数据文件。 |
229+| core*_instr_exe.csv| 代码指令详细信息，*代表0~n核，以便用户快速确定最耗时的指令，可参考代码指令信息文件。 |
230+| visualize_data.bin | 仿真流水图和仿真热点函数等信息可视化呈现文件。
231+|trace.json | 仿真指令流水图文件，包括每个核的子文件以及全部核的汇总文件，可参考指令流水图和内存通路吞吐率波形图。 |
181 232 
233+更多详细的信息可以参考[《算子开发工具用户指南》](https://hiascend.com/document/redirect/CannCommercialToolOpDev) 中的“算子调优 (msOpProf)”章节。