Star557
699
代码
代码
Issues49
Pull Requests174
Actions
项目讨论
Wiki
分析
Star557
699
已合并
[binary]上传Runtime编程指南，同时更新docs/README.md的文档大纲 #1030
ycm0028创建于 3月19日
[binary]上传Runtime编程指南，同时更新docs/README.md的文档大纲 #1030
已合并
ycm0028创建于 3月19日
从已删除 :master合入到cann/runtimemaster
共 712 个文件变更+3833-558
Adocs/01_quick_start/.gitattributes+1-0
@@ -0,0 +1 @@
1+*.png filter=lfs diff=lfs merge=lfs -text
Adocs/01_quick_start/Runtime简介.md+18-0
@@ -0,0 +1,18 @@
1+# CANN Runtime简介
2+ 
3+CANN Runtime是CANN软件栈中负责驱动硬件执行与管理AI计算任务的核心组件，它通过提供统一的API，使得上层应用和框架能够高效利用AI处理器的硬件计算资源。其主要功能包括以下几类
4+ 
5+- **Device管理**：提供对计算设备设置、重置、查询和配置等功能。
6+- **Memory管理**：提供设备内存和主机内存的申请、释放、内存拷贝等功能。
7+- **Context管理**： Context是计算设备的执行上下文环境，提供上下文创建、销毁，切换和配置等功能。
8+- **Stream管理**：Stream是Device提供的逻辑任务执行队列，提供Stream创建和销毁、属性查询和配置、Stream同步、Stream状态管理等功能。
9+- **Kernel管理**：提供AI Core \(AI Cube 、AI Vector\)、AI CPU等算子注册管理和KernelLaunch功能。
10+- **Event管理**：Event是用于设备内Stream间任务同步的事件，提供Event创建和销毁、事件同步、记录事件时间戳信息等功能。
11+- **Notify管理**：Notify主要是跨设备间的通知，提供Notify创建和销毁、Record/Wait功能。
12+- **运行时全局管理**：提供运行时初始化和进程级配置、DFX配置（例如算子数据Dump、溢出检测等）功能。
13+- **ACL Graph**：提供任务捕获或编排方式构建“运行时图”功能，利用“运行时图”任务整体下发能力，节省Host调度耗时提升整体性能。
14+ 
15+下图展示了Runtime的功能架构：
16+ 
17+![](figures/逻辑架构图.png)
18+ 
Adocs/01_quick_start/Runtime编程模型.md+262-0
@@ -0,0 +1,262 @@
1+# Runtime编程模型
2+ 
3+## 主机、设备编程模式
4+ 
5+**主机（或称Host）**：指X86服务器CPU、ARM服务器CPU，通过总线与一个或多个设备互联，利用设备提供的NN（Neural-Network）等计算能力完成业务。
6+ 
7+**设备（或称Device、NPU**）：指安装了AI处理器的硬件，通过总线（例如PCIe、HCCS等）与主机相连，提供NN等计算能力。HCCS是Huawei Cache Coherence System，表示华为缓存一致性系统。
8+ 
9+主机与设备的关系如下图所示：
10+ 
11+![](figures/主机设备关系图.png)
12+ 
13+总结如下：
14+ 
15+1. **主机和设备各自拥有独立的内存空间。**Runtime提供了主机侧和设备侧的内存申请接口，以及主机与设备之间的内存复制接口。在编程时，需要区分主机和设备的内存申请，并显式调用内存复制接口，将数据从主机侧复制到设备侧，以使设备硬件加速器在访问本地内存时达到最佳性能。
16+2. **主机与设备之间采用异步并行的执行方式。**
17+ 
18+ 主机将任务（或称Task）下发到设备后，**不会等待设备任务执行完成**就立即返回；设备随即开始调度并执行下发的任务；主机侧的CPU可以与设备侧的加速器并行工作。异步任务下发的接口通常会带有Stream参数，表示将任务下发到对应的Stream中执行。
19+ 
20+ 当主机需要获取设备的计算结果时，必须发起显式的同步API调用，同步API会阻塞主机CPU，直到设备侧任务执行完成才返回。
21+ 
22+ 通过这种异步并行执行机制，可以有效隐藏主机处理时间或主机与设备间的数据传输延迟，提高吞吐量，缩短端到端的执行时间。
23+ 
24+3. **异步任务下发到Stream中任务的执行方式**
25+ 
26+ Stream中的任务保序执行，Stream间的任务并行执行。例如下图中主机侧顺序启动Kernel1、Kernel2和Kernel3任务，具体执行顺序如下：
27+ 
28+ - Kernel1和Kernel3位于同一个Stream中，因此Kernel3需要等待Kernel1执行完毕才能开始执行。
29+ - Kernel2与Kernel1、Kernel3不在同一个Stream中，因此Kernel2可以与Kernel1、Kernel3并行执行。
30+ 
31+ ![](figures/Stream中的异步任务执行.png)
32+ 
33+## 典型执行流程
34+ 
35+基于Runtime编程的典型执行流程图如下所示：
36+ 
37+![](figures/典型执行流程图.png)
38+ 
39+1. Device初始化，以Device 0为例。
40+ 
41+ ```
42+ int32_t devId=0;
43+ aclrtSetDevice(devId);
44+ ```
45+ 
46+ 接口内部涉及如下操作：
47+ 1. 创建并初始化Device对象。
48+ 2. 为Device创建默认Context。
49+ 3. 为默认Context创建默认流。
50+ 4. 启动Device侧的CPU执行器进程。
51+ 
52+2. 在当前Context下创建Stream。
53+ 
54+ ```
55+ aclrtStream stream1;
56+ aclrtCreateStream(&stream1);
57+ ```
58+ 
59+ 接口内部涉及如下操作：
60+ 
61+ 1. 调用驱动创建任务队列。
62+ 2. Runtime侧创建Stream对象，Stream与任务队列关联。
63+ 3. 将Stream纳入到当前Context中管理。
64+ 
65+3. 申请主机内存。
66+ 
67+ ```
68+ uint64_t size=1024;
69+ void *hostPtr=nullptr;
70+ aclrtMallocHost(&hostPtr, size);
71+ ```
72+ 
73+4. 申请设备内存。
74+ 
75+ ```
76+ void *devPtr=nullptr;
77+ aclrtMalloc(&devPtr, size);
78+ ```
79+ 
80+5. Host到Device的内存同步拷贝。
81+ 
82+ ```
83+ aclrtMemcpy(devPtr, size, hostPtr, size, ACL_MEMCPY_HOST_TO_DEVICE);
84+ ```
85+ 
86+ Host到Device的内存同步拷贝，主要包括以下操作：
87+ 1. 根据源、目的地址构造DMA（Direct Memory Access）描述符后下发DMA任务。
88+ 2. 等待DMA任务完成，接口返回。
89+ 
90+6. 下发myKernel计算任务。
91+ 
92+ ```
93+ myKernel<<<numBlocks, nullptr, stream1>>>(devPtr);
94+ ```
95+ 
96+ 下发计算任务及任务调度的详细步骤如下：
97+ 
98+ 1. 向stream1下发myKernel计算任务，numBlocks用于指定计算任务的block数量，详细操作如下：
99+ 
100+ (a) myKernel的二进制代码注册并加载至Device （仅首次）。
101+ 
102+ (b) 将myKernel的执行参数拷贝至Device。
103+ 
104+ (c) 分配任务管理资源，并创建任务描述符。
105+ 
106+ (d) 下发任务到此流对应的任务队列。
107+ 
108+ (e) 下发完任务即接口异步返回。
109+ 
110+ 2. 任务调度器开始调度，详细操作如下：
111+ (a) 调度器按任务队列的优先级采用绝对优先级（SP）方式调度。
112+ 
113+ (b) 解析任务队列中下发的任务描述符。
114+ 
115+ (c) 根据任务类型判断对应加速单元是否有空闲。
116+ 
117+ (d) 有空闲且满足任务的核数量要求，则将任务调度给对应加速单元（例如，CPU算子会调度给CPU执行器）。
118+ 
119+ (e) 加速单元执行任务。
120+ 
121+ (f) 调度器等待加速单元执行完成，刷新任务状态和加速单元忙闲状态。当任务执行异常，调度器会处理异常。
122+ 
123+ (g) 继续调度其他任务。
124+ 
125+7. Stream同步。
126+ 
127+ ```
128+ aclrtSynchronizeStream(stream1);
129+ ```
130+ 
131+ Stream同步，详细操作如下：
132+ 1. 接口会阻塞当前CPU线程。
133+ 2. 通过轮循+中断机制同步流对应的任务队列执行状态。
134+ 3. 流上任务全部完成，此接口返回。
135+ 
136+8. Device到Host的内存同步拷贝，将结果复制回Host。
137+ 
138+ ```
139+ aclrtMemcpy(hostPtr, size, devPtr, size, ACL_MEMCPY_DEVICE_TO_HOST);
140+ ```
141+ 
142+ 根据源、目的地址构造DMA描述符后下发DMA任务。等待DMA任务完成，接口返回。
143+ 
144+9. 释放设备核主机内存。
145+ 
146+ ```
147+ aclrtFree(devPtr);
148+ aclrtFreeHost(hostPtr);
149+ ```
150+ 
151+10. 释放Device资源。
152+ 
153+ ```
154+ aclrtResetDeviceForce(devId);
155+ ```
156+ 
157+
158+ 
159+ 
160+## Runtime主要编程概念
161+ 
162+- **Host**, 是Runime对主机的抽象。
163+- **Device**，是对AI处理器所属设备的抽象，通常Host与Device关系为1：N。用户APP可以调用acl接口，例如aclrtSetDevice，指定当前用于运算的硬件设备。
164+- **Context**，是Device的逻辑运行环境，Context与Device的关系为N：1，即每个Context必定隶属于一个唯一的Device。Context负责管理运行资源对象（包括Stream、Event和Notify，但不包括内存）的生命周期；不同Context中的对象是完全隔离的，例如，不同Context的Stream和Event是完全隔离的，无法建立同步等待关系；运行出错同样按Context隔离。
165+- **Stream**，是Device提供的逻辑任务执行队列，可以异步地向Stream中添加任务，在同一个Stream中的任务会严格按FIFO方式执行。Stream与Context的关系是N：1，某条Stream一定属于唯一的Context。
166+- **Task**，可被添加到Stream中的执行任务，可以分计算类任务、内存拷贝、事件同步类任务。Task与Stream的关系是N：1，某个Task会被加入到唯一的Stream。
167+ 
168+Device、Context、Stream之间的关系如下图所示：
169+ 
170+![](figures/Device_Context_Stream关系.png)
171+ 
172+## 线程关联Context
173+ 
174+Runtime的大多数API接口没有device id参数，因为这些API接口所作用的Device是从调用线程关联的Context中获取的。因此，当主机线程调用Runtime API时，要遵循如下要求：
175+ 
176+- 线程（主机侧的CPU线程）要关联Context后，才能正确调用Runtime API。
177+- 线程同一时刻只能关联一个Context。
178+- 应用程序可以显式创建Context来达成运行资源隔离的业务诉求。此场景中，同一进程内Context可被所有线程可见，线程可以通过aclrtGetCurrentContext，aclrtSetCurrentContext进行切换Context。
179+ 
180+以下示例说明了线程在调用Context相关接口时，线程与Context之间的关联和切换过程，仅供参考，不可以直接拷贝编译运行。
181+ 
182+```
183+// 初始时，线程未关联任何Context
184+aclInit(nullptr);
185+aclrtSetDevice(0); // aclrtSetDevice会创建默认Context，同时将线程关联默认Context
186+// 线程关联的Context: 默认context
187+ 
188+aclrtContext ctx1, ctx2, current_ctx;
189+aclrtCreateContext(&ctx1, 0); // Device 0显式创建ctx1, 会将线程关联ctx1
190+// 线程关联的Context: ctx1
191+aclrtGetCurrentContext(&current_ctx); // 获取当前线程关联的Context, 此时返回的current_ctx==ctx1
192+// 线程关联的Context: ctx1
193+ 
194+aclrtCreateContext(&ctx2, 0); // Device 0又显式创建ctx2, 会将线程关联ctx2
195+// 线程关联的Context: ctx2
196+ 
197+aclrtSetCurrentContext(current_ctx); // 切换Context，由于current_ctx=ctx1，线程关联ctx1
198+// 线程关联的context: ctx1
199+ 
200+aclrtSetCurrentContext(ctx2); // 切换Context
201+// 线程关联的context: ctx2
202+.....
203+ 
204+aclrtDestroyContext(ctx2); // 当前线程正关联ctx2，销毁ctx2时会同时将线程也ctx2去关联
205+// 线程关联的context: NA
206+ 
207+// ctx2已销毁，应该切换到其他ctx(如ctx1)
208+aclrtSetCurrentContext(ctx1);
209+// 线程关联的context: ctx1
210+.....
211+ 
212+aclrtDestroyContext(ctx1); // 当前线程正关联ctx1，销毁ctx1时会同时将线程也ctx1去关联
213+// 线程关联的context: NA
214+aclrtResetDeviceForce(0);
215+```
216+ 
217+## 默认Context和默认Stream的使用场景
218+ 
219+- Device上执行操作下发前，必须有Context和Stream，这个Context、Stream可以显式创建，也可以隐式创建。**隐式创建**的Context、Stream就是默认Context、默认Stream。
220+ 
221+ 默认Stream作为接口入参时，直接传NULL。
222+ 
223+- **默认Context**不允许用户执行aclrtGetCurrentContext或aclrtSetCurrentContext操作，也不允许执行aclrtDestroyContext操作。
224+- **默认Context、默认Stream**一般适用于简单应用，用户仅需要一个Device的计算场景下。多线程应用程序建议使用显式创建的Context和Stream。
225+ 
226+示例代码如下，仅供参考，不可以直接拷贝编译运行：
227+ 
228+```
229+// ......
230+uint32_t numBlocks = 32;
231+uint64_t size = 1024;
232+void *devPtr = nullptr;
233+aclInit(nullptr);
234+aclrtSetDevice(0);
235+/* 已经创建了一个默认Context，在默认Context中创建了一个默认Stream，并且在当前线程可用 */
236+ 
237+......
238+aclrtMalloc(&devPtr, size );
239+myKernel<<<numBlocks, nullptr, nullptr>>>(devPtr); // <<< >>>中第三参数nullptr表示在默认Stream上执行
240+aclrtSynchronizeStream(nullptr);
241+ 
242+/* 等待计算任务全部完成，用户根据需要获取计算任务的输出结果 */
243+......
244+aclrtResetDeviceForce(0); // 释放Device 0，对应的默认Context及默认Stream生命周期也终止
245+```
246+ 
247+## 编写高性能应用程序的建议
248+ 
249+遵循如下基本原则：
250+ 
251+1. 主机侧与Device侧执行异步执行。主机侧要能及时下发足够任务至Device，确保加速硬件始终处于计算状态。
252+2. 采用多Stream方式充分利用Device上不同种类硬件加速器实现并发执行。 如下图所示，CANN Runtime可以协同调度多种硬件加速器，不同代AI处理器支持的硬件加速器不同，需以实际硬件用户手册中的说明为准。
253+ 
254+ ![](figures/多种硬件加速器.png)
255+ 
256+推荐如下方式：
257+ 
258+- 单线程中创建并使用多个Stream。如果单线程的性能足以满足向多个Stream下发任务，以充分利用Device的算力，建议采用单线程模式。
259+- 当单线程性能不足时，可以采用多线程模式来提升主机侧任务下发的性能。推荐每个线程创建并使用各自的Stream下发任务；不推荐多个线程并发向同一个Stream下发任务，这将引入锁操作，且多个线程间下发的任务是乱序的。
260+- Stream上下发的单个任务占不满AI Core时，可以使用多Stream下发可并行执行的任务来充分利用AI Core资源。
261+- AI处理器中包含多种硬件加速器，例如AI Core、AI CPU、DVPP（Digital Vision Pre-Processing）、Random（随机数生成器）等，这些硬件加速器对应不同类型的任务，建议多Stream的创建按照算子执行硬件划分。
262+ 
Adocs/01_quick_start/figures/Device_Context_Stream关系.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:853fedb3dc95e05eb9b1e56a038c43bff298a23d3ea03d6cae6fac9f900c79b4
3+size 45675
Adocs/01_quick_start/figures/Stream中的异步任务执行.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:14e21c3b3ac5732d02b3763814049136c585aff692d1e7863bcfb76e17d5f48d
3+size 36685
Adocs/01_quick_start/figures/主机设备关系图.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:50e67f34dc0280613e964552325eff81e2e1036dfa1a8432fd43643c258688b8
3+size 23859
Adocs/01_quick_start/figures/典型执行流程图.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:2f81bce209b760d299f29e1c127c20936289480e224e24da9d322fe419913ffb
3+size 57740
Adocs/01_quick_start/figures/多种硬件加速器.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:1dab0918f4653ab3ac8d9354f34bef65bbd59e6d911850ee74932ad2acce4cfa
3+size 33929
Adocs/01_quick_start/figures/逻辑架构图.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:49d774d4ad080d65a0d4b25787ae152c6c30c533c7f8c7b54a7eb5c801f7e362
3+size 34970
Adocs/01_quick_start/quick_start.md+5-0
@@ -0,0 +1,5 @@
1+# 快速入门
2+ 
3+- [Runtime简介](Runtime简介.md)
4+- [Runtime编程模型](Runtime编程模型.md)
5+ 
Adocs/02_dev_guide/.gitattributes+1-0
@@ -0,0 +1 @@
1+*.png filter=lfs diff=lfs merge=lfs -text
Adocs/02_dev_guide/ACL-Graph.md+8-0
@@ -0,0 +1,8 @@
1+# ACL Graph
2+ 
3+- **[单流捕获](单流捕获.md)**
4+ 
5+- **[跨流捕获](跨流捕获.md)**
6+ 
7+- **[任务更新](任务更新.md)**
8+ 
Adocs/02_dev_guide/Device内存使用.md+73-0
@@ -0,0 +1,73 @@
1+# Device内存使用
2+ 
3+在昇腾异构计算编程中，典型的使用场景是：通过aclrtMallocHost接口申请Host内存，通过aclrtMalloc接口申请Device内存，通过aclrtMemcpy（同步）/aclrtMemcpyAsync（异步）接口将数据从Host拷贝到Device上，算子执行过程中使用Device内存进行计算并保存结果。
4+ 
5+以下是一段简单的示例代码。在示例代码中，两个张量从Host内存被拷贝到Device内存，在Device侧完成计算，再将结果从Device内存拷贝到Host内存：
6+ 
7+```
8+int main(void)
9+{
10+ int32_t deviceId = 0;
11+ int64_t N = 16;
12+ const size_t bytes = static_cast<size_t>(N) * sizeof(float);
13+ aclrtStream stream = nullptr;
14+ 
15+ // STEP 1: 初始化、Stream创建
16+ aclInit(nullptr);
17+ aclrtSetDevice(deviceId);
18+ aclrtCreateStream(&stream);
19+ 
20+ // STEP 2: 申请Host内存
21+ void* hostA = nullptr;
22+ void* hostB = nullptr;
23+ void* hostOut = nullptr;
24+ 
25+ aclrtMallocHost(&hostA, bytes);
26+ aclrtMallocHost(&hostB, bytes);
27+ aclrtMallocHost(&hostOut, bytes);
28+ 
29+ // 输入数据初始化
30+ ...
31+ 
32+ // STEP 3: 申请Device内存
33+ void* deviceA = nullptr;
34+ void* deviceB = nullptr;
35+ void* deviceOut = nullptr;
36+ 
37+ // 第三个参数aclrtMemMallocPolicy表示申请内存时的内存页分配策略
38+ // ACL_MEM_MALLOC_HUGE_FIRST 表示大页内存优先，其余定义参考API文档
39+ aclrtMalloc(&deviceA, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
40+ aclrtMalloc(&deviceB, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
41+ aclrtMalloc(&deviceOut, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
42+ 
43+ // STEP 4: 将输入数据从Host内存传输到Device内存
44+ aclrtMemcpy(deviceA, bytes, hostA, bytes, ACL_MEMCPY_HOST_TO_DEVICE);
45+ aclrtMemcpy(deviceB, bytes, hostB, bytes, ACL_MEMCPY_HOST_TO_DEVICE);
46+ 
47+ // STEP 5: 执行计算操作，比如aclnnAdd，将计算结果保存在Device内存
48+ // ...
49+ 
50+ // STEP 6: 将计算结果从Device内存传输到Host内存
51+ aclrtMemcpy(hostOut, bytes, deviceOut, bytes, ACL_MEMCPY_DEVICE_TO_HOST);
52+ 
53+ // STEP 7: 在Host侧处理计算结果
54+ // ...
55+ 
56+ // STEP 8: 资源清理
57+ // STEP 8.1: 释放Host和Device内存
58+ if (deviceA) (void)aclrtFree(deviceA);
59+ if (deviceB) (void)aclrtFree(deviceB);
60+ if (deviceOut) (void)aclrtFree(deviceOut);
61+ if (hostA) (void)aclrtFreeHost(hostA);
62+ if (hostB) (void)aclrtFreeHost(hostB);
63+ if (hostOut) (void)aclrtFreeHost(hostOut);
64+ 
65+ // STEP 8.2: 清理设备和进程资源
66+ if (stream) (void)aclrtDestroyStream(stream);
67+ (void)aclrtResetDevice(deviceId);
68+ (void)aclFinalize();
69+ 
70+ return 0;
71+}
72+```
73+ 
Adocs/02_dev_guide/Event同步.md+27-0
@@ -0,0 +1,27 @@
1+# Event同步
2+ 
3+调用aclrtSynchronizeEvent接口阻塞当前主机线程直到指定的Event事件完成。以下为示例代码，不可以直接拷贝编译运行，仅供参考：
4+ 
5+```
6+// 创建Event
7+aclrtEvent event;
8+aclrtCreateEventExWithFlag(&event, ACL_EVENT_CAPTURE_STREAM_PROGRESS);
9+ 
10+// 创建Stream
11+aclrtStream stream;
12+aclrtCreateStream(&stream);
13+ 
14+// 在Stream上下发任务
15+......
16+ 
17+// 在Stream上记录Event
18+aclrtRecordEvent(event, stream);
19+ 
20+// 阻塞应用程序运行直到Event发生
21+aclrtSynchronizeEvent(event);
22+ 
23+// 显式销毁资源
24+aclrtDestroyStream(stream);
25+aclrtDestroyEvent(event);
26+```
27+ 
Adocs/02_dev_guide/Event查询.md+39-0
@@ -0,0 +1,39 @@
1+# Event查询
2+ 
3+调用aclrtQueryEventStatus接口可查询指定Event是否完成，非阻塞接口，Event状态包括ACL\_EVENT\_RECORDED\_STATUS\_COMPLETE（所有任务都已经执行完成）、ACL\_EVENT\_RECORDED\_STATUS\_NOT\_READY（有未执行完的任务）。
4+ 
5+以下是通过Event实现多线程内存池复用管理机制的代码示例，不可以直接拷贝编译运行，仅供参考。
6+ 
7+1. 在A线程中创建内存池，算子所用的内存来源于内存池，在算子后面插入Event Record任务。
8+ 
9+ ```
10+ // 申请内存池
11+ ......
12+ // 创建Stream
13+ aclrtStream stream;
14+ aclrtCreateStream(&stream);
15+
16+ // 创建Event
17+ aclrtEvent event;
18+ aclrtCreateEventExWithFlag(&event, ACL_EVENT_CAPTURE_STREAM_PROGRESS);
19+
20+ // 在Stream中下发计算任务
21+ ......
22+ // 在算子所在stream插入event
23+ aclrtRecordEvent(event, stream);
24+ ```
25+ 
26+2. 在B线程中调用查询接口，如果查询的Event已经完成，则代表Event Record前面的算子内存都可以被安全的复用。
27+ 
28+ ```
29+ aclrtEventRecordedStatus status;
30+ // 查询线程A的event是否完成
31+ aclrtQueryEventStatus(event, &status);
32+ if (status == ACL_EVENT_RECORDED_STATUS_COMPLETE) {
33+ // event已完成，算子占用的内存可以复用
34+ } else {
35+ // 算子并未执行完成，该算子占用的内存不能被复用
36+ }
37+
38+ ```
39+ 
Adocs/02_dev_guide/Event概念.md+12-0
@@ -0,0 +1,12 @@
1+# Event概念
2+ 
3+Event用于同一**Device内**、**不同Stream之间**的任务同步事件。**它支持一个任务等待一个事件**，例如stream2的任务依赖stream1的任务，想保证stream1中的任务先完成，这时可创建一个Event，将该Event插入到stream1中（Event Record任务），在stream2中插入一个等待Event完成的任务（Event Wait任务）；**也支持多个任务等待同一个事件（多等一）**，例如stream2和stream3中的任务都等待Stream1中的Event完成；同时，Event支持记录**事件时间戳**信息。
4+ 
5+一个任务等待一个事件的图示如下：
6+ 
7+![](figures/Event_一个任务等待一个事件.png)
8+ 
9+多个任务等待同一个事件的图示如下：
10+ 
11+![](figures/Event_多个任务等待一个事件.png)
12+ 
Adocs/02_dev_guide/Event的创建与销毁.md+16-0
@@ -0,0 +1,16 @@
1+# Event的创建与销毁
2+ 
3+以下是创建两个Event并销毁的代码示例，该示例仅用于说明Event使用方法，不可以直接拷贝编译运行。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/blob/master/example/event/1_event_timestamp)。
4+ 
5+```
6+aclrtEvent startEvent;
7+aclrtEvent endEvent;
8+// 创建Event，接口传入ACL_EVENT_SYNC参数，表示创建的Event用于同步
9+aclrtCreateEventExWithFlag(&startEvent, ACL_EVENT_SYNC);
10+aclrtCreateEventExWithFlag(&endEvent, ACL_EVENT_SYNC);
11+......
12+// 销毁Event
13+aclrtDestroyEvent(startEvent);
14+aclrtDestroyEvent(endEvent);
15+```
16+ 
Adocs/02_dev_guide/Event等待.md+45-0
@@ -0,0 +1,45 @@
1+# Event等待
2+ 
3+多Stream之间任务的同步等待可以利用Event实现，例如，若stream2的任务依赖stream1的任务，想保证stream1中的任务先完成，这时可创建一个Event，调用aclrtRecordEvent接口将Event插入到stream1中（通常称为Event Record任务），调用aclrtStreamWaitEvent接口在stream2中插入一个等待Event完成的任务（通常称为Event Wait任务）。
4+ 
5+以下为调用aclrtStreamWaitEvent接口的示例代码，不可以直接拷贝编译运行，仅供参考：
6+ 
7+```
8+// 创建一个Event
9+aclrtEvent event;
10+aclrtCreateEventExWithFlag(&event, ACL_EVENT_SYNC);
11+ 
12+// 创建stream1
13+aclrtStream stream1;
14+aclrtCreateStream(&stream1);
15+ 
16+// 创建stream2
17+aclrtStream stream2;
18+aclrtCreateStream(&stream2);
19+ 
20+// 在stream1上下发任务
21+......
22+ 
23+// 在stream1末尾添加了一个event
24+aclrtRecordEvent(event, stream1);
25+ 
26+// 在stream2上下发不依赖stream1执行完成的任务
27+......
28+ 
29+// 阻塞stream2运行，直到指定event发生，也就是stream1执行完成
30+aclrtStreamWaitEvent(stream2, event);
31+ 
32+// 在stream2上下发依赖stream1执行完成的任务
33+......
34+ 
35+// 阻塞应用程序运行，直到stream1和stream2中的所有任务都执行完成
36+aclrtSynchronizeStream(stream1);
37+aclrtSynchronizeStream(stream2);
38+ 
39+// 显式销毁资源
40+aclrtDestroyStream(stream1);
41+aclrtDestroyStream(stream2);
42+aclrtDestroyEvent(event);
43+......
44+```
45+ 
Adocs/02_dev_guide/Event管理.md+14-0
@@ -0,0 +1,14 @@
1+# Event管理
2+ 
3+- **[Event概念](Event概念.md)**
4+ 
5+- **[Event的创建与销毁](Event的创建与销毁.md)**
6+ 
7+- **[Event等待](Event等待.md)**
8+ 
9+- **[记录Event时间戳](记录Event时间戳.md)**
10+ 
11+- **[Event查询](Event查询.md)**
12+ 
13+- **[Event同步](Event同步.md)**
14+ 
Adocs/02_dev_guide/Host回调任务.md+37-0
@@ -0,0 +1,37 @@
1+# Host回调任务
2+ 
3+CANN为CPU和NPU之间的异步协作提供了灵活的方式。用户可以使用aclrtLaunchHostFunc在Stream的任意位置插入一个Host回调任务。当本Stream上所有前序任务执行完成后，该Host回调任务会被自动执行，并且会阻塞本Stream上的后续任务执行。
4+ 
5+回调函数不能直接或者间接调用CANN Runtime API，否则可能会导致错误或死锁。
6+ 
7+以下是在Stream上插入一个Host回调任务的代码示例，不可以直接拷贝编译运行，仅供参考。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/tree/master/example/callback/1_callback_hostfunc)。
8+ 
9+```
10+// Host回调任务
11+void myHostCallback(void *args)
12+{
13+ printf("In MyHostCallback.\n");
14+ 
15+ // myKernel1完成后的处理，阻塞MyKernel2的执行
16+ ......
17+}
18+......
19+ 
20+// 创建Stream
21+aclrtStream stream;
22+aclrtCreateStream(&stream);
23+ 
24+// 在Stream上下发任务
25+aclrtMemcpyAsync(devPtrIn, size, hostPtr, hostSize, ACL_MEMCPY_HOST_TO_DEVICE, stream);
26+myKernel1<<<8, nullptr, stream>>>(devPtrIn, devPtrOut, size);
27+aclrtLaunchHostFunc(stream, myHostCallback, nullptr);
28+myKernel2<<<8, nullptr, stream>>>(devPtrOut, size);
29+aclrtMemcpyAsync(hostPtr, hostSize, devPtrOut, size, ACL_MEMCPY_DEVICE_TO_HOST, stream);
30+ 
31+// 阻塞应用程序运行，直到指定Stream中的所有任务都完成
32+aclrtSynchronizeStream(stream);
33+ 
34+// 销毁Stream
35+aclrtDestroyStream(stream);
36+```
37+ 
Adocs/02_dev_guide/Host锁页内存使用.md+68-0
@@ -0,0 +1,68 @@
1+# Host锁页内存使用
2+ 
3+在CANN编程框架中，Host内存可以是**Pageable内存**，也可以是**Page-Locked内存**（也称锁页内存或Pinned内存）：
4+ 
5+- **Pageable内存**，由操作系统统一管理。开发者可使用malloc、mmap等传统接口申请内存，使用free、munmap等接口释放内存**。**当内存压力较大时，Pageable内存会被换出到后备存储提供的交换空间中。当Pageable内存中的数据被传输到Device时，数据首先会被复制到缓冲区，然后通过DMA（Direct Memory Access）通道传输到Device。
6+- **Page-Locked内存，**即锁页内存。开发者需要用Runtime提供的API进行锁页内存的申请和释放，例如aclrtMallocHost、aclrtFreeHost等接口**。**对于锁页内存，虚拟页与物理页的映射关系固定，在其生命周期内不会被换出至交换空间。当锁页内存中的数据被传输至Device时，直接通过DMA通道传输，无需经过缓冲区，传输性能更优。
7+ 
8+ 在Runtime中，**使用锁页内存的好处如下**：
9+ 
10+ - 设备可以直接通过DMA访问主机内存，无需经过缓冲区，可以提供更好的传输性能；
11+ - 数据搬运过程无需CPU参与，可以实现数据的异步传输；
12+ - 数据的异步传输，使传输过程和计算过程可以相互掩盖，减少传输+计算的整体时长。
13+ 
14+锁页内存可直接通过aclrtMallocHost接口申请，示例代码如下。如果需要在申请时指定内存的其他配置，例如自定义模块ID、配置VA（virtual address）一致性等，也可以使用aclrtMallocHostWithCfg接口申请。
15+ 
16+```
17+// 资源初始化
18+......
19+ 
20+// 申请锁页内存
21+void *hostPtr = NULL;
22+aclrtMallocHost(&hostPtr, size);
23+ 
24+// 申请Device内存
25+void *devicePtr = NULL;
26+aclrtMalloc(&devicePtr, size, ACL_MEM_MALLOC_NORMAL_ONLY);
27+ 
28+// 内存初始化
29+......
30+ 
31+// 异步H2D
32+aclrtMemcpyAsync(devicePtr, size, hostPtr, size, ACL_MEMCPY_HOST_TO_DEVICE, stream);
33+ 
34+// 等待异步拷贝完成
35+aclrtSynchronizeStream(stream);
36+ 
37+// 资源释放
38+if (devicePtr) (void)aclrtFree(devicePtr);
39+if (hostPtr) (void)aclrtFreeHost(hostPtr);
40+......
41+```
42+ 
43+若使用malloc/mmap等内存管理接口申请Pageable内存，当前Runtime提供了aclrtHostRegisterV2接口，用于将Pageable内存转换为锁页内存，供Device访问。请注意，当OS内核版本为5.10或更低时，此方法会导致异常，此时应通过aclrtMallocHost接口申请锁页内存。
44+ 
45+```
46+// 申请Pageable内存
47+void *hostPtr = malloc(size);
48+ 
49+// 注册为锁页内存(ACL_HOST_REG_PINNED)，并映射到Device(ACL_HOST_REG_MAPPED)
50+aclrtHostRegisterV2(hostPtr, size, ACL_HOST_REG_PINNED|ACL_HOST_REG_MAPPED);
51+ 
52+// 获取Device地址
53+void *devicePtr = NULL;
54+aclrtHostGetDevicePointer(hostPtr, &devicePtr, 0);
55+ 
56+// 任务通过Device地址访问对应内存
57+...
58+ 
59+// 资源释放
60+aclrtHostUnregister(hostPtr);
61+free(hostPtr);
62+```
63+ 
64+若涉及Host内存的VA（virtual address）一致性：
65+ 
66+- Host内存可以通过aclrtHostRegister接口注册到Device，根据Host指针映射得到Device指针，供Device使用；默认情况下，对于同一块Host内存，Host和Device看到的虚拟地址是不同的；
67+- 如果需要Host和Device的虚拟地址保持一致，可以使用aclrtMallocHostWithCfg接口申请锁页内存，并在aclrtMallocConfig参数中指定ACL\_RT\_MEM\_ATTR\_VA\_FLAG属性，后续再注册获取到的Device虚拟地址即可与Host虚拟地址一致。
68+ 
Adocs/02_dev_guide/Kernel加载与执行.md+118-0
@@ -0,0 +1,118 @@
1+# Kernel加载与执行
2+ 
3+Kernel函数可以采用<<<\>\>\>方式进行任务下发，具有代码简洁，可读性好的优点。
4+ 
5+以下是关键步骤的代码示例，不可以直接拷贝编译运行，仅供参考。
6+ 
7+```
8+// Device code
9+Template<>
10+extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
11+{
12+ KernelAdd op;
13+ op.Init(x, y, z);
14+ op.Process();
15+}
16+ 
17+int main()
18+{
19+ int N = ...;
20+ size_t size = N * sizeof(uint64);
21+ 
22+ // Initialize
23+ aclrtSetDevice(0);
24+ 
25+ // Create stream
26+ aclrtStream stream;
27+ aclrtCreateStream(&stream);
28+ 
29+ // Allocate vectors in host memory
30+ void *h_x, *h_y, *h_z;
31+ aclrtMallocHost(&h_x, size);
32+ aclrtMallocHost(&h_y, size);
33+ aclrtMallocHost(&h_z, size);
34+ 
35+ // Initialize input vectors
36+ ...
37+ 
38+ // Allocate vectors in device memory
39+ void *d_x, *d_y, *d_z;
40+ aclrtMalloc(&d_x, size, ACL_MEM_MALLOC_HUGE_FIRST);
41+ aclrtMalloc(&d_y, size, ACL_MEM_MALLOC_HUGE_FIRST);
42+ aclrtMalloc(&d_z, size, ACL_MEM_MALLOC_HUGE_FIRST);
43+ 
44+ // Copy vectors from host memory to device memory
45+ aclrtMemcpy(d_x, size, h_x, size, ACL_MEMCPY_HOST_TO_DEVICE);
46+ aclrtMemcpy(d_y, size, h_y, size, ACL_MEMCPY_HOST_TO_DEVICE);
47+ 
48+ // Invoke kernel
49+ uint32_t numBlocks = 48;
50+ add_custom<<<numBlocks, nullptr, stream>>>(d_x, d_y, d_z);
51+ ...
52+}
53+```
54+ 
55+用户也可以使用Runtime提供的LaunchKernel接口（例如aclrtLaunchKernelWithHostArgs接口）进行kernel函数的下发。使用这种方式需先了解Binary和Function的概念：
56+ 
57+- Binary：是一个动态加载的代码容器单元，里面包含编译后的kernel代码、全局变量等。用户可以通过aclrtBinaryLoadFromFile或aclrtBinaryLoadFromData将编译好的算子二进制加载到NPU上，并获得对应的Binary句柄。
58+- Function：是一个具体可执行的kernel函数，它定义在Binary内部，是主机代码可以调用并在NPU上执行的入口点。用户可以通过aclrtBinaryGetFunction获取kernel函数对应的Function句柄。
59+ 
60+以下是使用LaunchKernel接口的关键代码示例，不可以直接拷贝编译运行，仅供参考。
61+ 
62+```
63+// Device code
64+extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
65+{
66+ KernelAdd op;
67+ op.Init(x, y, z);
68+ op.Process();
69+}
70+ 
71+int main()
72+{
73+ int N = ...;
74+ size_t size = N * sizeof(uint64);
75+ 
76+ // Initialize
77+ aclrtSetDevice(0);
78+ 
79+ // Create stream
80+ aclrtStream stream;
81+ aclrtCreateStream(&stream);
82+ 
83+ // Create binary from binary file
84+ aclrtBinHandle bin;
85+ aclrtBinaryLoadFromFile("add_custom.o", nullptr, &bin);
86+ 
87+ // Get function handle from binary
88+ aclrtFuncHandle add_custom;
89+ aclrtBinaryGetFunction(bin, "add_custom", &add_custom);
90+ 
91+ // Allocate vectors in host memory
92+ void *h_x, *h_y, *h_z;
93+ aclrtMallocHost(&h_x, size);
94+ aclrtMallocHost(&h_y, size);
95+ aclrtMallocHost(&h_z, size);
96+ 
97+ // Initialize input vectors
98+ ...
99+ 
100+ // Allocate vectors in device memory
101+ void *d_x, *d_y, *d_z;
102+ aclrtMalloc(&d_x, size, ACL_MEM_MALLOC_HUGE_FIRST);
103+ aclrtMalloc(&d_y, size, ACL_MEM_MALLOC_HUGE_FIRST);
104+ aclrtMalloc(&d_z, size, ACL_MEM_MALLOC_HUGE_FIRST);
105+ 
106+ // Copy vectors from host memory to device memory
107+ aclrtMemcpy(d_x, size, h_x, size, ACL_MEMCPY_HOST_TO_DEVICE);
108+ aclrtMemcpy(d_y, size, h_y, size, ACL_MEMCPY_HOST_TO_DEVICE);
109+ 
110+ // Invoke kernel
111+ uint32_t numBlocks = 48;
112+ void* args[] = {d_x, d_y, d_z};
113+ size_t argsSize = 3 * sizeof(void*);
114+ aclrtLaunchKernelWithHostArgs(add_custom, numBlocks, stream, nullptr, args, argsSize, nullptr, 0);
115+ ...
116+}
117+```
118+ 
Adocs/02_dev_guide/Notify管理.md+44-0
@@ -0,0 +1,44 @@
1+# Notify管理
2+ 
3+Notify通常用于多个Device之间的同步，如下图所示：Device 0向Device 1发送完数据后，通过Notify通知Device 1数据已写完。
4+ 
5+![](figures/Notify_多个Device之间的同步.png)
6+ 
7+Notify只支持一对一通知机制。若要实现向多个Device发起通知，要发起多次Notify操作，如下图所示：
8+ 
9+![](figures/Notify_多个Device通知机制.png)
10+ 
11+Notify与Event功能区别在于，Notify Wait完成后，Notify状态会自动重置，因此一个Notify Record任务只能通知一个Notify Wait任务；而Event Wait并不会自动重置Event状态，因此一个Event Record任务可以做到通知一个或多个Event Wait任务。此外，Notify不支持时间戳功能。
12+ 
13+在同一个Device两条流间同步的场景下，Notify可以实现Event实现相同的效果。
14+ 
15+Notify相关接口的调用代码示例如下：
16+ 
17+```
18+// 创建Stream
19+aclrtStream stream1;
20+aclrtStream stream2;
21+aclrtCreateStream(&stream1);
22+aclrtCreateStream(&stream2);
23+// 创建Notify
24+aclrtNotify notify;
25+aclrtCreateNotify(&notify, ACL_NOTIFY_DEFAULT);
26+ 
27+// 在stream2插入wait
28+aclrtWaitAndResetNotify(notify, stream2, 0);
29+// 在stream2中下发计算任务
30+......
31+ 
32+// 在stream1下发stream2计算任务所依赖的拷贝任务
33+......
34+// 在stream1插入record
35+aclrtRecordNotify(notify, stream1);
36+ 
37+// 同步等待任务完成
38+aclrtSynchronizeStream(stream1);
39+aclrtSynchronizeStream(stream2);
40+ 
41+// 销毁notify
42+aclrtDestroyNotify(notify);
43+```
44+ 
Adocs/02_dev_guide/Persistent流.md+42-0
@@ -0,0 +1,42 @@
1+# Persistent流
2+ 
3+非Persistent流上的任务在执行完成之后从Stream出队。如果要多次执行某个任务，需要在非Persistent流上多次下发该任务。
4+ 
5+Runtime提供了Persistent流支持任务的持久化。在Persistent流上下发的任务不会被立即执行，任务执行完成后也不会被立即销毁。只有在销毁Persistent流时，相关的任务才会被销毁。
6+ 
7+调用aclrtCreateStreamWithConfig接口创建Persistent流，Persistent流需要与模型运行实例创建绑定，支持模型的反复执行。以下为示例代码，不可以直接拷贝编译运行，仅供参考：
8+ 
9+```
10+// 创建Persistent stream
11+aclrtStream stream;
12+aclrtCreateStreamWithConfig(&stream, 0, ACL_STREAM_PERSISTENT);
13+ 
14+// 构建一个模型运行实例
15+aclmdlRI modelRI;
16+aclmdlRIBuildBegin(&modelRI, 0);
17+ 
18+// 把Persistent stream绑定到模型运行实例
19+aclmdlRIBindStream(modelRI, stream, ACL_MODEL_STREAM_FLAG_HEAD);
20+ 
21+// 在Persistent流上下发任务
22+......
23+ 
24+// 标记下发任务结束
25+aclmdlRIEndTask(modelRI, stream);
26+ 
27+// 结束模型运行实例构建
28+aclmdlRIBuildEnd(modelRI, nullptr);
29+ 
30+// 在默认stream多次执行模型运行实例
31+aclmdlRIExecute(modelRI, -1);
32+aclmdlRIExecute(modelRI, -1);
33+ 
34+// 解除绑定
35+aclmdlRIUnbindStream(modelRI, stream);
36+ 
37+// 销毁资源
38+aclrtDestroyStream(stream);
39+aclmdlRIDestroy(modelRI);
40+......
41+```
42+ 
Adocs/02_dev_guide/Stream创建与销毁.md+20-0
@@ -0,0 +1,20 @@
1+# Stream创建与销毁
2+ 
3+调用aclrtCreateStream创建Stream，得到的aclrtStream对象作为后续的内存异步复制、Stream同步、Kernel执行等接口的Stream入参。显式创建的Stream需要调用aclrtDestroyStream接口显式销毁。销毁Stream时，如果Stream上有未完成的任务，则会等待任务完成后再销毁Stream。
4+ 
5+以下是创建Stream并在Stream上下发计算任务的代码示例，不可以直接拷贝编译运行，仅供参考。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/blob/master/example/stream/0_simple_stream)。
6+ 
7+```
8+// 显式创建一个Stream
9+aclrtStream stream;
10+aclrtCreateStream(&stream);
11+ 
12+// 在Stream上下发Host->Device复制任务、MyKernel任务、和Device->Host复制任务
13+aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize, ACL_MEMCPY_HOST_TO_DEVICE, stream);
14+myKernel<<<8, nullptr, stream>>>();
15+aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize, ACL_MEMCPY_DEVICE_TO_HOST, stream);
16+ 
17+// 销毁Stream（等待Device->Host复制任务执行完成后销毁）
18+aclrtDestroyStream(stream);
19+```
20+ 
Adocs/02_dev_guide/Stream和Event的行为.md+24-0
@@ -0,0 +1,24 @@
1+# Stream和Event的行为
2+ 
3+在与当前Device无所属关系的Stream上下发算子将会失败，示例代码如下：
4+ 
5+```
6+aclrtSetDevice(0); // 指定Device 0作为计算设备
7+aclrtStream s0;
8+aclrtCreateStream(&s0); // 在Device 0上创建Stream s0
9+myKernel<<<8, nullptr, s0>>>(); // 在Device 0上通过Stream s0下发算子
10+ 
11+aclrtSetDevice(1); // 指定Device 1作为计算设备
12+aclrtStream s1;
13+aclrtCreateStream(&s1); // 在Device 1上创建Stream s1
14+myKernel<<<8, nullptr, s1>>>(); // 在Device 1上通过Stream s1下发算子
15+ 
16+// 算子下发失败
17+myKernel<<<8, nullptr, s0>>>(); // 在Device 1上通过Stream s0下发算子
18+```
19+ 
20+- 当Stream所属的Device和当前操作的Device不相同时，在此Stream上调用aclrtMemcpyAsync会失败。
21+- 当Event和Stream关联到不同的Device上时，调用aclrtRecordEvent会失败。
22+- 当Event和Stream关联到不同的Device上时，调用aclrtStreamWaitEvent会失败。
23+- 当Event所属的Device和当前操作的Device不相同时，aclrtSynchronizeEvent和aclrtQueryEvent会成功。
24+ 
Adocs/02_dev_guide/Stream概念.md+12-0
@@ -0,0 +1,12 @@
1+# Stream概念
2+ 
3+Stream描述了一个在Host下发并在Device上执行的任务队列。
4+ 
5+在同一个Stream中，任务按照进入队列的顺序依次执行。当硬件资源充足时，不同Stream上的任务会被调度到不同的硬件资源上并行执行。当硬件资源不足时，不同Stream上的任务可能串行执行。
6+ 
7+Stream可以配置优先级、遇错即停、persistent等多种属性。优先级会影响不同Stream上的任务的执行顺序。通常情况下，高优先级Stream上的任务会优先于低优先级Stream上的任务执行。如果Stream配置了persistent属性，则下发在其上的任务不会被立即执行，执行完成后也不会被立即销毁。Persistent Stream适用于模型运行实例构建场景。
8+ 
9+相对于主机线程，Stream上的任务是异步执行的。主机线程可以使用Device同步接口来等待当前Context下所有Stream上的任务全部执行完成，或者使用Stream同步接口来等待Stream上的任务全部执行完成。
10+ 
11+Runtime中的Stream均为非阻塞式Stream，默认Stream不会跟显式创建的Stream进行隐式同步。
12+ 
Adocs/02_dev_guide/Stream管理.md+18-0
@@ -0,0 +1,18 @@
1+# Stream管理
2+ 
3+- **[Stream概念](Stream概念.md)**
4+ 
5+- **[Stream创建与销毁](Stream创建与销毁.md)**
6+ 
7+- **[默认Stream](默认Stream.md)**
8+ 
9+- **[显式同步](显式同步.md)**
10+ 
11+- **[Host回调任务](Host回调任务.md)**
12+ 
13+- **[配置Stream优先级](配置Stream优先级.md)**
14+ 
15+- **[配置任务遇错即停](配置任务遇错即停.md)**
16+ 
17+- **[Persistent流](Persistent流.md)**
18+ 
Adocs/02_dev_guide/dev_guide.md+51-0
@@ -0,0 +1,51 @@
1+# Runtime编程指南
2+ 
3+- [初始化](初始化.md)
4+- [内存管理](内存管理.md)
5+ - [内存管理总述](内存管理总述.md)
6+ - [Device内存使用](Device内存使用.md)
7+ - [Host锁页内存使用](Host锁页内存使用.md)
8+ - [虚拟内存管理](虚拟内存管理.md)
9+ 
10+- [异步任务执行](异步任务执行.md)
11+ - [异步任务总述](异步任务总述.md)
12+ - [Stream管理](Stream管理.md)
13+ - [Stream概念](Stream概念.md)
14+ - [Stream创建与销毁](Stream创建与销毁.md)
15+ - [默认Stream](默认Stream.md)
16+ - [显式同步](显式同步.md)
17+ - [Host回调任务](Host回调任务.md)
18+ - [配置Stream优先级](配置Stream优先级.md)
19+ - [配置任务遇错即停](配置任务遇错即停.md)
20+ - [Persistent流](Persistent流.md)
21+ 
22+ - [Kernel加载与执行](Kernel加载与执行.md)
23+ - [Event管理](Event管理.md)
24+ - [Event概念](Event概念.md)
25+ - [Event的创建与销毁](Event的创建与销毁.md)
26+ - [Event等待](Event等待.md)
27+ - [记录Event时间戳](记录Event时间戳.md)
28+ - [Event查询](Event查询.md)
29+ - [Event同步](Event同步.md)
30+ 
31+ - [Notify管理](Notify管理.md)
32+ - [内存语义同步](内存语义同步.md)
33+ - [系统任务](系统任务.md)
34+ 
35+- [ACL Graph](ACL-Graph.md)
36+ - [单流捕获](单流捕获.md)
37+ - [跨流捕获](跨流捕获.md)
38+ - [任务更新](任务更新.md)
39+ 
40+- [多设备编程](多设备编程.md)
41+ - [多设备编程总述](多设备编程总述.md)
42+ - [多Device选择](多Device选择.md)
43+ - [Stream和Event的行为](Stream和Event的行为.md)
44+ - [跨Device的数据交互](跨Device的数据交互.md)
45+ 
46+- [进程间通信](进程间通信.md)
47+- [运行时核资源控制](运行时核资源控制.md)
48+- [配置AI Core栈空间大小](配置AI-Core栈空间大小.md)
49+- [异常处理](异常处理.md)
50+ - [获取Runtime错误码](获取Runtime错误码.md)
51+ 
Adocs/02_dev_guide/figures/ACL_Graph任务更新流程1.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:a1f15443b52b4079b4e2a9035b3edb7b1d2ed32ac6fc8993525f27aa23f535fd
3+size 54385
Adocs/02_dev_guide/figures/ACL_Graph任务更新流程2.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:efdb880071f956daf789040b9ad2771bc98dc0623c2f718c057ad33c3a467cb3
3+size 69330
Adocs/02_dev_guide/figures/ACL_Graph任务更新流程3.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:d2e25b087af2c92ec1512cdf014637270731fc423938ac8720682f7f7ddc09ef
3+size 99318
Adocs/02_dev_guide/figures/ACL_Graph单流捕获.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:6337102d2bc2f45037831d7b01801344d25c06bed29d9e105635177a2df3809f
3+size 52407
Adocs/02_dev_guide/figures/ACL_Graph跨流捕获.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:95a51a3c47fc825b35b46d687f5d45126aae16a3e8d81d591ed3fd98e26be980
3+size 70407
Adocs/02_dev_guide/figures/Event_一个任务等待一个事件.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:bebcde165c394697340978128f262219f4c2f619e4fa20a2a0122fea9505fc52
3+size 20116
Adocs/02_dev_guide/figures/Event_多个任务等待一个事件.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:dd64f7a6a9a3130a7ea3cc5fc69565b2473daf9f392c262443a6e5e06ff48033
3+size 31836
Adocs/02_dev_guide/figures/Notify_多个Device之间的同步.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:dbae9383edc2ce1b3aef044626f439af9122c87006cc8cbce6041efa38562b05
3+size 33908
Adocs/02_dev_guide/figures/Notify_多个Device通知机制.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:1a6cc323896ef8aeb5bc0f42810acf05aa3b232ed4c0cef8ed5dc73f9046395b
3+size 22594
Adocs/02_dev_guide/figures/内存语义同步.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:b060a85c998c60088b6bda59e5b296a46cf1908191a5d6a4c0e544edc821131d
3+size 34823
Adocs/02_dev_guide/figures/同步等待流程_多Device场景.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:3b2ad03111fe95c154720c9351c3cd1cb5a21271dcc749f5cdaadf3bfc88e2b4
3+size 168127
Adocs/02_dev_guide/figures/进程间共享内存.png+3-0
@@ -0,0 +1,3 @@
1+version https://git-lfs.github.com/spec/v1
2+oid sha256:7c895ba89eadc9a10db8c313f87be6ab6acd365eee1bd1d636d066e66fa62b31
3+size 72039
Adocs/02_dev_guide/任务更新.md+232-0
@@ -0,0 +1,232 @@
1+# 任务更新
2+ 
3+**须知：**本功能为试验特性，后续版本可能会存在变更，不支持应用于商用产品中。
4+当Stream上的任务已经被捕获，并暂存到模型中之后，若要更新任务（包含任务本身以及任务的参数信息），当前支持两种方式：
5+ 
6+1. **方式一**：以需要更新的任务为分界点，在aclmdlRICaptureBegin和aclmdlRICaptureEnd接口之间分别捕获该任务前后的任务、并暂存到不同的模型中，分开执行。
7+ 
8+ **该方式适用于大量任务需要更新的场景**（例如一个模型有两种不同Shape的input输入场景）**，接口调用逻辑比较简单，但导致暂存捕获任务的模型数量增多，若模型数量超出硬件资源限制，则会触发报错。**
9+ 
10+ **该方式的基本使用流程如下图所示：**
11+ 
12+ ![](figures/ACL_Graph任务更新流程1.png)
13+ 
14+2. **方式二**：在aclmdlRICaptureBegin、aclmdlRICaptureEnd接口之间下发主流上需捕获的任务，通过aclmdlRICaptureTaskGrpBegin、aclmdlRICaptureTaskGrpEnd接口将待更新的任务标记为在一个任务组中，并返回任务组的handle，在aclmdlRICaptureTaskUpdateBegin、aclmdlRICaptureTaskUpdateEnd接口之间更新任务。
15+ 
16+ **该方式适用于少量单算子调用任务需要更新的场景，支持先更新任务再依次执行模型实例中的任务，也支持更新任务与模型实例中其他任务的并发执行。但更新任务比单独下发任务更耗时，另外，还存在一些使用限制**：aclmdlRICaptureTaskGrpBegin、aclmdlRICaptureTaskGrpEnd接口之间的任务数量、任务类型，要与aclmdlRICaptureTaskUpdateBegin、aclmdlRICaptureTaskUpdateEnd接口之间的任务数量、任务类型一致；跨Stream捕获任务的场景下，在aclmdlRICaptureTaskGrpBegin、aclmdlRICaptureTaskGrpEnd接口之间，其它捕获状态的Stream上不允许同时下发任务；任务组类似一个临界资源，不支持多线程多Stream并发更新，否则会导致更新结果非预期。
17+ 
18+ - **对于“先更新任务，再依次执行aclmdlRI实例中的任务”的场景，使用流程如下图所示：**
19+ 
20+ ![](figures/ACL_Graph任务更新流程2.png)
21+ 
22+ - **对于“更新任务与其他任务的并发执行”的场景，使用流程如下图所示：**
23+ 
24+ 若模型的运行实例中存在大量任务，为了提升性能，可使用external类型的Event实现更新任务与其他任务的并发执行，并且需再单独创建一个用于更新任务的Stream（下文称之为UpdateStream）。这里的external类型的Event，是指调用aclrtCreateEventWithFlag接口并设置flag为ACL\_EVENT\_EXTERNAL的Event，该类型的Event规格有限，且无法实现跨Stream的任务捕获，需要考虑合理复用。创建external类型的Event之后，在UpdateStream上下发更新任务，接着调用aclrtRecordEvent接口下发一个Event Record任务。然后，在主流中，在待更新的任务之前，调用aclrtStreamWaitEvent接口下发一个Event Wait任务，用于等待UpdateStream中的任务更新完成。最后，在主流中，调用aclrtStreamWaitEvent接口之后，再调用aclrtResetEvent接口重置external类型的Event。
25+ 
26+ ![](figures/ACL_Graph任务更新流程3.png)
27+ 
28+ 以任务并发执行的场景为例，以下是更新aclnnAdd算子输入参数的关键代码示例。
29+ 
30+ ```
31+ #include <stdio.h>
32+ #include <vector>
33+ #include "acl/acl.h"
34+ #include "aclnnop/aclnn_add.h"
35+
36+ #define ACL_LOG(fmt, args...) fprintf(stdout, "[INFO] " fmt "\n", ##args)
37+
38+ int64_t GetShapeSize(const std::vector<int64_t> &shape)
39+ {
40+ int64_t shape_size = 1;
41+ for (auto i : shape) {
42+ shape_size *= i;
43+ }
44+ return shape_size;
45+ }
46+
47+ int CreateAclTensor(const std::vector<int64_t> &shape, void **deviceAddr,
48+ aclDataType dataType, aclTensor **tensor)
49+ {
50+ auto size = GetShapeSize(shape) * sizeof(float);
51+ // 申请Device侧内存
52+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
53+ // 计算连续tensor的stride
54+ std::vector<int64_t> strides(shape.size(), 1);
55+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
56+ strides[i] = shape[i + 1] * strides[i + 1];
57+ }
58+ // 调用aclCreateTensor接口创建aclTensor
59+ *tensor = aclCreateTensor(shape.data(),
60+ shape.size(),
61+ dataType,
62+ strides.data(),
63+ 0,
64+ aclFormat::ACL_FORMAT_ND,
65+ shape.data(),
66+ shape.size(),
67+ *deviceAddr);
68+ return 0;
69+ }
70+
71+ int main()
72+ {
73+ int devID = 0;
74+ void *self_d = nullptr;
75+ void *other_d = nullptr;
76+ void *out_d = nullptr;
77+ void *outtmp_d = nullptr;
78+ aclTensor *self = nullptr;
79+ aclTensor *other = nullptr;
80+ aclScalar *alpha = nullptr;
81+ aclScalar *updatealpha = nullptr;
82+ aclTensor *out = nullptr;
83+ aclTensor *outtmp = nullptr;
84+ /* aclnnAdd: self = self + other * alpha */
85+ float *self_h = nullptr;
86+ float *other_h = nullptr;
87+ std::vector<int64_t> shape = {4, 2};
88+ float *out_h = nullptr;
89+ float alphaValue = 1.1f;
90+ float updatealphaValue = 5.5f;
91+ uint64_t workspaceSize = 0;
92+ uint64_t workspaceSize1 = 0;
93+ uint64_t workspaceSize2 = 0;
94+ aclOpExecutor *executor2;
95+ aclOpExecutor *executor;
96+ aclOpExecutor *executor1;
97+ auto size = GetShapeSize(shape);
98+
99+ // 初始化
100+ aclInit(NULL);
101+ // 指定计算设备
102+ aclrtSetDevice(devID);
103+
104+ // 准备aclnnAdd算子的输入、输出参数
105+ CreateAclTensor(shape, &self_d, aclDataType::ACL_FLOAT, &self);
106+ CreateAclTensor(shape, &other_d, aclDataType::ACL_FLOAT, &other);
107+ alpha = aclCreateScalar(&alphaValue, aclDataType::ACL_FLOAT);
108+ updatealpha = aclCreateScalar(&updatealphaValue, aclDataType::ACL_FLOAT);
109+ CreateAclTensor(shape, &out_d, aclDataType::ACL_FLOAT, &out);
110+ CreateAclTensor(shape, &outtmp_d, aclDataType::ACL_FLOAT, &outtmp);
111+
112+ // 调用aclnnAdd算子的第一段接口，获取算子计算所需的workspace大小以及包含了算子计算流程的执行器
113+ // 后续涉及多次调用aclnnAdd算子，此处需调用多次第一段接口，获取不同的aclOpExecutor
114+ // outtmp = self + alpha * other
115+ // 更新前：out = outtmp + alpha * other 更新后：out = outtmp + updatealpha * other
116+ aclnnAddGetWorkspaceSize(self, other, alpha, outtmp, &workspaceSize, &executor);
117+ void *workspaceAddr = nullptr;
118+ if (workspaceSize > 0) {
119+ aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
120+ }
121+ // 更新前：out = outtmp + alpha * other
122+ aclnnAddGetWorkspaceSize(outtmp, other, alpha, out, &workspaceSize1, &executor1);
123+ void *workspaceAddr1 = nullptr;
124+ if (workspaceSize1 > 0) {
125+ aclrtMalloc(&workspaceAddr1, workspaceSize1, ACL_MEM_MALLOC_HUGE_FIRST);
126+ }
127+ // 更新后：out = outtmp + updatealpha * other
128+ aclnnAddGetWorkspaceSize(outtmp, other, updatealpha, out, &workspaceSize2, &executor2);
129+ void *workspaceAddr2 = nullptr;
130+ if (workspaceSize2 > 0) {
131+ aclrtMalloc(&workspaceAddr2, workspaceSize2, ACL_MEM_MALLOC_HUGE_FIRST);
132+ }
133+
134+ // 使用aclrtMallocHost申请锁页内存
135+ aclrtMallocHost((void **)&self_h, size * sizeof(float));
136+ aclrtMallocHost((void **)&other_h, size * sizeof(float));
137+ aclrtMallocHost((void **)&out_h, size * sizeof(float));
138+ for (int i = 0; i < 8; i++) {
139+ self_h[i] = static_cast<float>(0);
140+ other_h[i] = static_cast<float>(1);
141+ out_h[i] = static_cast<float>(0);
142+ }
143+
144+ aclmdlRI modelRI;
145+ aclrtStream stream1;
146+ aclrtCreateStream(&stream1);
147+ aclrtEvent event;
148+
149+ // 创建external类型的event
150+ aclrtCreateEventWithFlag(&event, ACL_EVENT_EXTERNAL);
151+
152+ // ========开始捕获任务========
153+ aclmdlRICaptureBegin(stream1, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL);
154+ // 异步拷贝，将aclnnAdd算子self输入的数据从Host侧传到Device侧
155+ aclrtMemcpyAsync(self_d, size * sizeof(float), self_h, size * sizeof(float), ACL_MEMCPY_HOST_TO_DEVICE, stream1);
156+ // 异步拷贝，将aclnnAdd算子other输入的数据从Host侧传到Device侧
157+ aclrtMemcpyAsync(other_d, size * sizeof(float), other_h, size * sizeof(float), ACL_MEMCPY_HOST_TO_DEVICE, stream1);
158+ // 执行aclnnAdd算子
159+ aclnnAdd(workspaceAddr, workspaceSize, executor, stream1);
160+ // 在主流stream1上，下发一个Event Wait任务，等待更新任务完成
161+ aclrtStreamWaitEvent(stream1, event);
162+ aclrtResetEvent(event, stream1);
163+ aclrtTaskGrp handle;
164+ // 标记要更新的任务
165+ aclmdlRICaptureTaskGrpBegin(stream1);
166+ aclnnAdd(workspaceAddr1, workspaceSize1, executor1, stream1);
167+ aclmdlRICaptureTaskGrpEnd(stream1, &handle);
168+ // 异步拷贝，将算子输出数据从Device侧传回Host侧
169+ aclrtMemcpyAsync(out_h, size * sizeof(float), out_d, size * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST, stream1);
170+ // ========结束捕获任务========
171+ aclmdlRICaptureEnd(stream1, &modelRI);
172+
173+ aclrtStream updateStream;
174+ aclrtCreateStream(&updateStream);
175+
176+ for (int i = 0; i < 2; i++) {
177+ ACL_LOG("execute model, loop: %d", i);
178+ aclmdlRIExecuteAsync(modelRI, stream1);
179+ // 开始更新任务，将aclnnAdd算子的alpha参数更新为updatealpha
180+ aclmdlRICaptureTaskUpdateBegin(updateStream, handle);
181+ if (i == 1) {
182+ aclnnAdd(workspaceAddr2, workspaceSize2, executor2, updateStream);
183+ ACL_LOG("update alpha value of aclnnAdd");
184+ }
185+ aclmdlRICaptureTaskUpdateEnd(updateStream);
186+ // 更新任务之后，在updateStream上，下发Event Record任务，用于通知主流stream1继续执行Event Wait之后的任务
187+ aclrtRecordEvent(event, updateStream);
188+ aclrtSynchronizeStream(updateStream);
189+ aclrtSynchronizeStream(stream1);
190+ ACL_LOG("%f %f %f %f %f %f %f %f\n",
191+ out_h[0],
192+ out_h[1],
193+ out_h[2],
194+ out_h[3],
195+ out_h[4],
196+ out_h[5],
197+ out_h[6],
198+ out_h[7]);
199+ }
200+
201+ // 释放资源
202+ aclmdlRIDestroy(modelRI);
203+ aclDestroyTensor(self);
204+ aclDestroyTensor(other);
205+ aclDestroyTensor(out);
206+ aclDestroyTensor(outtmp);
207+ aclDestroyScalar(alpha);
208+ aclDestroyScalar(updatealpha);
209+ aclrtFree(self_d);
210+ aclrtFree(other_d);
211+ aclrtFree(out_d);
212+ aclrtFree(outtmp_d);
213+ aclrtDestroyStream(stream1);
214+ aclrtDestroyStream(updateStream);
215+ aclrtDestroyEvent(event);
216+ if (workspaceAddr != nullptr) {
217+ aclrtFree(workspaceAddr);
218+ }
219+ if (workspaceAddr1 != nullptr) {
220+ aclrtFree(workspaceAddr1);
221+ }
222+ if (workspaceAddr2 != nullptr) {
223+ aclrtFree(workspaceAddr2);
224+ }
225+ // 释放计算设备的资源
226+ aclrtResetDevice(devID);
227+ // 去初始化
228+ aclFinalize();
229+ }
230+
231+ ```
232+ 
Adocs/02_dev_guide/内存管理.md+10-0
@@ -0,0 +1,10 @@
1+# 内存管理
2+ 
3+- **[内存管理总述](内存管理总述.md)**
4+ 
5+- **[Device内存使用](Device内存使用.md)**
6+ 
7+- **[Host锁页内存使用](Host锁页内存使用.md)**
8+ 
9+- **[虚拟内存管理](虚拟内存管理.md)**
10+ 
Adocs/02_dev_guide/内存管理总述.md+9-0
@@ -0,0 +1,9 @@
1+# 内存管理总述
2+ 
3+在昇腾异构计算架构中，系统由主机（Host）和设备（Device）组成。Host和Device各自拥有独立的内存，Host内存是指AI处理器所在服务器的主机内存（即CPU内存），而Device内存则是指AI处理器自带的设备内存。
4+ 
5+内存管理中要做好的两件事是：
6+ 
7+1. **可以访问内存：**Runtime提供了一套内存管理API，使开发者能够高效便捷地编写应用程序中的内存管理代码。由于Host和Device的内存相互独立，Runtime提供了专门的接口来分别申请和释放Host内存及Device内存。例如，申请和释放Host内存的接口为aclrtMallocHost和aclrtFreeHost，而申请和释放Device内存的接口为aclrtMalloc和aclrtFree。
8+2. **高效访问内存：**为了实现最佳的内存访问性能，需要将数据存储在对应的内存中，例如算子在Device上执行过程中，访问Device上的数据性能要远高于访问Host侧的。为此，Runtime提供了Host与Device之间互相拷贝内存的接口，支持同步和异步方式，例如aclrtMemcpy和aclrtMemcpyAsync等，以便开发者更好地规划数据的存储与访问。
9+ 
Adocs/02_dev_guide/内存语义同步.md+64-0
@@ -0,0 +1,64 @@
1+# 内存语义同步
2+ 
3+内存语义同步机制允许用户基于通用Device内存实现同步。与Event/Notify同步机制不同，基于内存语义的同步机制还支持算子作为同步参与方，即算子可以在执行过程中与另一条流进行同步。
4+ 
5+下图展示了在算子执行过程中与另一条流之间的同步过程：
6+ 
7+![](figures/内存语义同步.png)
8+ 
9+内存语义同步相关接口的调用代码示例如下，不可以直接拷贝编译运行，仅供参考：
10+ 
11+- **Device示例代码**（算子核函数实现代码）
12+ 
13+ ```
14+ extern "C" __global__ __aicore__ void myKernel1(GM_ADDR syncMem)
15+ {
16+ // 算子逻辑
17+ ......
18+ // 向syncMem所指向内存写1
19+ __gm__ uint64_t* flag = reinterpret_cast<__gm__ uint64_t*>(syncMem);
20+ *flag = 1;
21+ dcci(flag, 0, 2);
22+ // 算子逻辑
23+ ......
24+ }
25+ extern "C" __global__ __aicore__ void myKernel2(GM_ADDR syncMem)
26+ {
27+ // 算子逻辑
28+ ...
29+ // 轮询阻塞直到syncMem所指向内存值为2
30+ __gm__ volatile uint64_t* flag = reinterpret_cast<__gm__ uint64_t*>(syncMem);
31+ dcci(flag, 0, 2);
32+ while (*flag != 2) {
33+ dcci(flag, 0, 2);
34+ }
35+ // 算子逻辑
36+ ......
37+ }
38+ ```
39+ 
40+- **Host示例代码**
41+ 
42+ ```
43+ // 创建Stream
44+ aclrtStream stream1;
45+ aclrtStream stream2;
46+ aclrtCreateStream(&stream1);
47+ aclrtCreateStream(&stream2);
48+
49+ // 申请Device内存
50+ void* syncMem;
51+ aclrtMalloc(&syncMem, sizeof(uint64_t), ACL_MEM_MALLOC_NORMAL_ONLY);
52+ // 在stream1上下发wait任务，该任务阻塞等待直到syncMem所指向内存中的值为1
53+ aclrtValueWait(syncMem, 1, ACL_VALUE_WAIT_EQ, stream1);
54+ // 在stream2上下发myKernel1，该kernel内部向syncMem所指向内存写1，从而解除stream1上wait任务的阻塞状态
55+ myKernel1<<<numBlocks, nullptr, stream2>>>(syncMem);
56+
57+ // 在stream1上下发myKernel2，该kernel内部轮询等待直到syncMem所指向内存的值为2
58+ myKernel2<<<numBlocks, nullptr, stream1>>>(syncMem);
59+ // 在stream2上下发write任务，该任务往syncMem所指向内存写为2，从而解除stream1上myKernel2的阻塞状态
60+ aclrtValueWrite(syncMem, 2, 0, stream2);
61+ ```
62+ 
63+**说明：**因为内存语义同步机制是基于通用Device内存实现，所以可以通过aclrtMemset/aclrtMemsetAsync初始化和清除同步所用的内存。
64+ 
Adocs/02_dev_guide/初始化.md+55-0
@@ -0,0 +1,55 @@
1+# 初始化
2+ 
3+**CANN Runtime**提供了aclInit、aclrtSetDevice接口，在应用程序启动时被调用，结合json配置文件完成如下功能：
4+ 
5+- **初始化环境**：设置 CANN Runtime 运行时所需的环境参数，确保所有运行时资源和配置项都被正确加载。
6+- **设备资源配置**：初始化与硬件相关的资源（如 Ascend 处理器、加速卡等）并为其分配资源，使得后续的计算任务可以在适当的设备上执行。
7+- **设置日志**：提供日志记录的初始化，确保系统的运行状态可以被实时检查与调试。
8+- **资源管理初始化**：为后续的内存管理、任务调度、内存分配等功能提供资源准备。
9+ 
10+以下是初始化及指定计算设备的代码示例，不可以直接拷贝编译运行，仅供参考。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/tree/master/example/device/0_device_normal)。
11+ 
12+```
13+// 初始化
14+int32_t deviceId = 0;
15+aclInit(nullptr); // json配置路径为nullptr, 默认初始化
16+aclrtSetDevice(deviceId);
17+ 
18+// SetDevice后才可以调用其他aclrt运行时接口。
19+......
20+ 
21+// 去初始化
22+aclrtResetDeviceForce(deviceId);
23+aclFinalize();
24+```
25+ 
26+若不显式调用aclrtSetDevice接口时，可按照如下配置的默认Device进行处理。例如：aclInit接口的json文件中指定默认Device 为 0：
27+ 
28+```
29+{
30+ "defaultDevice":{
31+ "default_device":"0"
32+ }
33+}
34+```
35+ 
36+在aclInit接口中启用默认Device功能后，初始化的示例代码如下，不可以直接拷贝编译运行，仅供参考。完整样例代码请参见[Link](https://gitcode.com/cann/runtime-dev/tree/master/example/device/0_device_normal)。
37+ 
38+```
39+// 初始化
40+int32_t deviceId = 0;
41+aclInit(nullptr);
42+ 
43+// 启用DefaultDevice后，可以不显式调用aclrtSetDevice直接调用运行时接口
44+// 接口中会按json配置文件指定的device，进行隐式aclrtSetDevice
45+aclrtMalloc(&devPtr, size, 0);
46+ 
47+......
48+ 
49+// 去初始化
50+aclrtResetDeviceForce(deviceId)
51+aclFinalize();
52+```
53+ 
54+除了默认Device功能之外，通过配置aclInit接口中的json文件，可实现零修改代码二进制完成性能数据采集、模型输入/输出数据导出、溢出算子Dump等功能，相关配置示例及使用说明请参考aclInit接口中的说明。
55+ 
Adocs/02_dev_guide/单流捕获.md+173-0
@@ -0,0 +1,173 @@
1+# 单流捕获
2+ 
3+**须知：**本功能为试验特性，后续版本可能会存在变更，不支持应用于商用产品中。
4+ 
5+在当前主流框架（例如PyTorch）采用的Eager模式下，每个操作或任务都是边下发边执行，无需构建计算图，这种模式带来了即时可见的执行效果和便捷的调试功能，但同时也带来了Host的下发开销。随着性能优化的不断深入，这些Host开销逐渐成为瓶颈，变成不可忽视的问题。
6+ 
7+在AI处理器上可以将相关任务下沉到Device上并执行，从而减少Host的开销。为了达到此效果，**当前提供了“捕获Stream任务到模型中、再执行模型”的acl接口**，如下图所示，在aclmdlRICaptureBegin和aclmdlRICaptureEnd接口之间，所有在指定Stream上下发的任务不会立即执行，而是被暂存在模型的运行实例中，只有在调用aclmdlRIExecuteAsync接口执行模型时这些任务才会被真正执行。当Stream上的任务需要被多次执行时，无需再下发任务，只需多次调用aclmdlRIExecuteAsync接口执行模型即可，达到减少Host侧的任务下发开销的效果。任务执行完毕后，若无需再使用模型的运行实例，可调用aclmdlRIDestroy接口及时销毁该资源。
8+ 
9+捕获任务到模型中再执行模型的基本流程如下图所示：
10+ 
11+![](figures/ACL_Graph单流捕获.png)
12+ 
13+捕获任务到模型中、再执行模型的场景下，存在如下基本限制：
14+ 
15+1. 在进入捕获状态前，Stream上的任务依然是立即执行的。
16+2. 在Stream上捕获任务时，只会将任务下沉到Device上，并不会立即执行，因此，对Stream或Event的查询或同步均为非法操作。同样，对Device或Context的查询或同步也是非法的，因为Device和Context中包含了Stream的同步信息。捕获过程中，对Stream、Event、Device、Context的同步或查询，在任何捕获模式下都是非法的。
17+3. 在捕获过程中，在ACL\_MODEL\_RI\_CAPTURE\_MODE\_GLOBAL模式（全局禁止，所有线程都不可以调用非安全函数）下，调用内存同步操作类函数（例如aclrtMemset、aclrtMemcpy、aclrtMemcpy2d）是非法的，会校验报错导致捕获失败。若业务侧确定这些函数的执行不会影响任务捕获，此时，可以通过调用aclmdlRICaptureThreadExchangeMode接口切换当前线程的捕获模式为ACL\_MODEL\_RI\_CAPTURE\_MODE\_RELAXED，解除调用限制。
18+ 
19+4. 在捕获过程中，下发配置类的任务，例如Profiling配置、Dump配置、溢出检测配置等，可能会返回报错或者对捕获模型不生效。
20+5. 若捕获的异步内存复制任务涉及Host内存，则只支持使用acl接口（例如aclrtMallocHost）申请Host锁页内存，否则在捕获过程中将返回报错。
21+6. 另外，在捕获过程中，对默认Stream的操作是非法的。
22+7. 最后还需要注意的是，任务被捕获后，需要使用者保证模型中任务使用资源的有效性，直至模型被销毁后才能销毁相关资源。
23+ 
24+以下是单流捕获add算子计算的示例代码。
25+ 
26+```
27+#include <stdio.h>
28+#include <vector>
29+#include "acl/acl.h"
30+#include "aclnnop/aclnn_add.h"
31+ 
32+#define ACL_LOG(fmt, args...) fprintf(stdout, "[INFO] " fmt "\n", ##args)
33+ 
34+int64_t GetShapeSize(const std::vector<int64_t> &shape)
35+{
36+ int64_t shape_size = 1;
37+ for (auto i : shape) {
38+ shape_size *= i;
39+ }
40+ return shape_size;
41+}
42+ 
43+int CreateAclTensor(const std::vector<int64_t> &shape, void **deviceAddr,
44+ aclDataType dataType, aclTensor **tensor)
45+{
46+ auto size = GetShapeSize(shape) * sizeof(float);
47+ // 申请Device侧内存
48+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
49+ // 计算连续tensor的stride
50+ std::vector<int64_t> strides(shape.size(), 1);
51+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
52+ strides[i] = shape[i + 1] * strides[i + 1];
53+ }
54+ // 调用aclCreateTensor接口创建aclTensor
55+ *tensor = aclCreateTensor(shape.data(),
56+ shape.size(),
57+ dataType,
58+ strides.data(),
59+ 0,
60+ aclFormat::ACL_FORMAT_ND,
61+ shape.data(),
62+ shape.size(),
63+ *deviceAddr);
64+ return 0;
65+}
66+ 
67+int main()
68+{
69+ int devID = 0;
70+ void *self_d = nullptr;
71+ void *other_d = nullptr;
72+ void *out_d = nullptr;
73+ aclTensor *self = nullptr;
74+ aclTensor *other = nullptr;
75+ aclScalar *alpha = nullptr;
76+ aclTensor *out = nullptr;
77+ /* aclnnAdd: out = self + other * alpha */
78+ float *self_h = nullptr;
79+ float *other_h = nullptr;
80+ std::vector<int64_t> shape = {4, 2};
81+ float alphaValue = 1.1f;
82+ uint64_t workspaceSize = 0;
83+ aclOpExecutor *executor;
84+ auto size = GetShapeSize(shape);
85+
86+ // 初始化
87+ aclInit(NULL);
88+ // 指定计算设备
89+ aclrtSetDevice(devID);
90+ 
91+ // 准备aclnnAdd算子的输入、输出参数
92+ CreateAclTensor(shape, &self_d, aclDataType::ACL_FLOAT, &self);
93+ CreateAclTensor(shape, &other_d, aclDataType::ACL_FLOAT, &other);
94+ alpha = aclCreateScalar(&alphaValue, aclDataType::ACL_FLOAT);
95+ CreateAclTensor(shape, &out_d, aclDataType::ACL_FLOAT, &out);
96+ 
97+ // 获取算子计算所需的workspace大小以及包含了算子计算流程的执行器
98+ aclnnAddGetWorkspaceSize(self, other, alpha, out, &workspaceSize, &executor);
99+ void *workspaceAddr = nullptr;
100+ if (workspaceSize > 0) {
101+ aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
102+ }
103+
104+ // 使用aclrtMallocHost申请锁页内存
105+ aclrtMallocHost((void **)&self_h, size * sizeof(float));
106+ aclrtMallocHost((void **)&other_h, size * sizeof(float));
107+ for (int i = 0; i < 8; i++) {
108+ self_h[i] = static_cast<float>(0);
109+ other_h[i] = static_cast<float>(1);
110+ }
111+ 
112+ aclmdlRI modelRI;
113+ aclrtStream stream;
114+ aclrtCreateStream(&stream);
115+ 
116+ // ========开始捕获任务========
117+ aclmdlRICaptureBegin(stream, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL);
118+ // 异步拷贝，将算子self输入的数据从Host侧传到Device侧
119+ aclrtMemcpyAsync(self_d, size * sizeof(float), self_h, size * sizeof(float), ACL_MEMCPY_HOST_TO_DEVICE, stream);
120+ // 切换捕获模式为RELAXED，允许调用aclrtMemcpy函数
121+ aclmdlRICaptureMode mode = ACL_MODEL_RI_CAPTURE_MODE_RELAXED;
122+ aclmdlRICaptureThreadExchangeMode(&mode);
123+ // 同步拷贝，将算子other输入的数据从Host侧传到Device侧，仅执行一次
124+ aclrtMemcpy(other_d, size * sizeof(float), other_h, size * sizeof(float), ACL_MEMCPY_HOST_TO_DEVICE);
125+ // 将捕获模式切换回GLOBAL
126+ aclmdlRICaptureThreadExchangeMode(&mode);
127+ // 执行aclnnAdd算子
128+ aclnnAdd(workspaceAddr, workspaceSize, executor, stream);
129+ // 异步拷贝，将算子输出数据从Device侧传回Host侧
130+ aclrtMemcpyAsync(self_h, size * sizeof(float), out_d, size * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST, stream);
131+ // ========结束捕获任务========
132+ aclmdlRICaptureEnd(stream, &modelRI);
133+
134+ // 打印模型信息，维测场景下使用
135+ const char *jsonPath = "./modelRI.json";
136+ aclmdlRIDebugJsonPrint(modelRI, jsonPath, 0);
137+ 
138+ // 多次执行模型
139+ for (int i = 0; i < 8; i++) {
140+ aclmdlRIExecuteAsync(modelRI, stream);
141+ aclrtSynchronizeStream(stream);
142+ // 打印每一次的算子输出数据
143+ ACL_LOG("%f %f %f %f %f %f %f %f\n",
144+ self_h[0],
145+ self_h[1],
146+ self_h[2],
147+ self_h[3],
148+ self_h[4],
149+ self_h[5],
150+ self_h[6],
151+ self_h[7]);
152+ }
153+ 
154+ // 释放资源
155+ aclmdlRIDestroy(modelRI);
156+ aclrtDestroyStream(stream);
157+ aclDestroyTensor(self);
158+ aclDestroyTensor(other);
159+ aclDestroyTensor(out);
160+ aclDestroyScalar(alpha);
161+ aclrtFree(self_d);
162+ aclrtFree(other_d);
163+ aclrtFree(out_d);
164+ if (workspaceAddr != nullptr) {
165+ aclrtFree(workspaceAddr);
166+ }
167+ // 释放计算设备的资源
168+ aclrtResetDevice(devID);
169+ // 去初始化
170+ aclFinalize();
171+}
172+```
173+ 
Adocs/02_dev_guide/多Device选择.md+50-0
@@ -0,0 +1,50 @@
1+# 多Device选择
2+ 
3+一个Host搭配多Device的场景下，用户可以在Host侧应用程序中通过aclrtGetDeviceCount接口来获取当前Host上搭配的Device数量，Device按照0、1、2、... 的顺序排布。
4+ 
5+以下为获取Device信息的代码示例，不可以直接拷贝编译运行，仅供参考：
6+ 
7+```
8+// 获取Device数量及其对应的属性信息
9+uint32_t deviceCount;
10+aclrtGetDeviceCount(&deviceCount);
11+uint32_t deviceId;
12+for (deviceId = 0; deviceId < deviceCount; ++deviceId) {
13+ // 按需查询设备属性信息
14+ aclrtDevAttr attr = ACL_DEV_ATTR_VECTOR_CORE_NUM;
15+ int64_t value;
16+ aclrtGetDeviceInfo(deviceId, attr, &value);
17+}
18+```
19+ 
20+此时，可以随时通过aclrtSetDevice接口按**线程粒度**切换Device（不会影响其他线程）。指定Device后，后续的内存分配、Kernel执行等操作均在该Device上进行，且Stream、Event等也与当前指定的Device相关联。用户可以通过aclrtResetDevice接口释放资源，但更推荐使用aclrtResetDeviceForce接口一次性清理Device上的资源，包括默认Context、默认Stream以及在默认Context下创建的所有Stream。如果默认Context或默认Stream下的任务尚未完成，系统会等待任务完成后才释放资源。在用户程序中，若使用aclrtResetDevice接口，则需确保aclrtSetDevice和aclrtResetDevice接口的调用次数成对出现。
21+ 
22+多Device选择的接口调用流程如下图所示：
23+ 
24+![](figures/同步等待流程_多Device场景.png)
25+ 
26+以下是多Device选择的代码示例，不可以直接拷贝编译运行，仅供参考：
27+ 
28+```
29+// 指定Device 0作为计算设备，并将Device 0的默认Context作为当前线程的默认Context
30+aclrtSetDevice(0);
31+aclrtStream s0;
32+aclrtCreateStream(&s0);
33+// 执行任务1
34+......
35+ 
36+// 指定Device 1作为计算设备，并将Device1的默认Context作为当前线程的默认Context
37+aclrtSetDevice(1);
38+aclrtStream s1;
39+aclrtCreateStream(&s1);
40+// 执行任务2
41+......
42+ 
43+// 复位Device 1，释放计算资源，线程默认Context被释放
44+// 如需进行继续运行任务，需要显示指定device&context
45+aclrtResetDeviceForce(1);
46+ 
47+// 复位device0，释放计算资源
48+aclrtResetDeviceForce(0);
49+```
50+ 
Adocs/02_dev_guide/多设备编程.md+10-0
@@ -0,0 +1,10 @@
1+# 多设备编程
2+ 
3+- **[多设备编程总述](多设备编程总述.md)**
4+ 
5+- **[多Device选择](多Device选择.md)**
6+ 
7+- **[Stream和Event的行为](Stream和Event的行为.md)**
8+ 
9+- **[跨Device的数据交互](跨Device的数据交互.md)**
10+ 
Adocs/02_dev_guide/多设备编程总述.md+11-0
@@ -0,0 +1,11 @@
1+# 多设备编程总述
2+ 
3+多设备编程使用户能够利用多个NPU（Neural-Network Processing Unit）的综合性能和内存等资源，实现超越单个NPU的性能水平。通常，每个物理NPU在Runtime编程界面中被抽象为一个Device，任务下发时需要有Device中的Context进行支撑，因此多设备编程需要管理多个Device及其相应的Context。
4+ 
5+一些常见的多NPU编程方法包括：
6+ 
7+- 单个主机线程驱动多个NPU。
8+- 多个主机线程，每个线程驱动自己的NPU。
9+- 多个单线程主机进程，每个进程驱动自己的NPU。
10+- 多个主机进程，每个进程包含多个线程，每个线程驱动自己的NPU。
11+ 
Adocs/02_dev_guide/异常处理.md+4-0
@@ -0,0 +1,4 @@
1+# 异常处理
2+ 
3+- **[获取Runtime错误码](获取Runtime错误码.md)**
4+ 
Adocs/02_dev_guide/异步任务总述.md+38-0
@@ -0,0 +1,38 @@
1+# 异步任务总述
2+ 
3+Runtime将下列操作视为彼此可并发执行的独立任务：
4+ 
5+- 主机计算
6+- 设备计算
7+- 主机到设备的内存复制
8+- 设备到主机的内存复制
9+- 设备内部内存复制
10+- 设备间内存复制
11+ 
12+这些操作之间的实际并发程度取决于设备的特性集和计算能力，详见下文：
13+ 
14+1. 主机与设备并发执行
15+ 
16+ 通过异步接口实现主机并发：接口在设备完成任务前就返回控制权给主机线程。利用异步调用，可以一次性将多个设备操作排队，在设备资源可用时执行，从而减轻主机线程的管理负担，使其能够执行其他任务。以下设备操作相对于主机是异步的：
17+ 
18+ - Kernel执行
19+ - 模型运行实例执行
20+ - 接口名带Async的计算任务
21+ - 接口名带Async的内存拷贝，涉及非锁页主机内存的异步内存拷贝也可能退化为同步。
22+ 
23+2. 多个Kernel并发执行
24+ 
25+ 同一个Context中下发的多个Kernel是可以并发执行的，不同Context间的多个Kernel同样可以并发执行。任务调度器会根据Stream的优先级和AI Core忙闲状态进行统一调度。
26+ 
27+3. 不同加速器任务并发执行
28+ 
29+ 除AI Core外，Device上还提供其他硬件加速器，例如AI CPU、DVPP（Digital Vision Pre-Processing）、Random（随机数生成器）等，不同的加速器任务之间并发执行。不同代AI处理器支持的硬件加速器不同，需以实际硬件用户手册中的说明为准。
30+ 
31+4. 数据传输与计算任务并发
32+ 
33+ 数据传输有独立的DMA（Direct Memory Access）引擎，可以与计算任务并发执行。
34+ 
35+5. 双向数据传输
36+ 
37+ 从Host到Device和从Device到Host的数据传输是可并发的。
38+ 
Adocs/02_dev_guide/异步任务执行.md+16-0
@@ -0,0 +1,16 @@
1+# 异步任务执行
2+ 
3+- **[异步任务总述](异步任务总述.md)**
4+ 
5+- **[Stream管理](Stream管理.md)**
6+ 
7+- **[Kernel加载与执行](Kernel加载与执行.md)**
8+ 
9+- **[Event管理](Event管理.md)**
10+ 
11+- **[Notify管理](Notify管理.md)**
12+ 
13+- **[内存语义同步](内存语义同步.md)**
14+ 
15+- **[系统任务](系统任务.md)**
16+ 
Adocs/02_dev_guide/显式同步.md+52-0
@@ -0,0 +1,52 @@
1+# 显式同步
2+ 
3+对于异步任务接口，主机线程调用异步任务接口后仅代表下发任务，不代表任务执行完成。用户需要显式调用设备同步、流同步等显式同步接口等待任务完成。调用此类显式同步接口后，主机线程会阻塞直到相关的任务执行完成。
4+ 
5+## 设备同步：aclrtSynchronizeDevice
6+ 
7+阻塞当前主机线程直到当前Device的当前Context中所有显式或隐式创建的Stream完成已下发的所有任务。
8+ 
9+以下是设备同步代码示例，不可以直接拷贝编译运行，仅供参考。
10+ 
11+```
12+// 指定Device
13+aclrtSetDevice(0);
14+ 
15+// 创建Stream
16+aclrtStream stream;
17+aclrtCreateStream(&stream);
18+ 
19+// 在Stream上下发任务
20+......
21+ 
22+// 阻塞应用程序运行，直到正在运算中的Device完成运算
23+aclrtSynchronizeDevice();
24+ 
25+// 资源销毁
26+aclrtDestroyStream(stream);
27+aclrtResetDevice(0);
28+```
29+ 
30+## 流同步：aclrtSynchronizeStream
31+ 
32+阻塞当前主机线程直到指定的Stream完成已下发的所有任务。
33+ 
34+以下是流同步的代码示例，不可以直接拷贝编译运行，仅供参考。
35+ 
36+```
37+// 创建Stream
38+aclrtStream stream;
39+aclrtCreateStream(&stream);
40+ 
41+// 在Stream上下发任务
42+......
43+ 
44+// 调用aclrtSynchronizeStream接口，阻塞应用程序运行，直到指定Stream中的所有任务都完成。
45+aclrtSynchronizeStream(stream);
46+ 
47+// Stream使用结束后，显式销毁Stream
48+aclrtDestroyStream(stream);
49+```
50+ 
51+此外，用户可以使用aclrtStreamQuery查询stream上的任务是否全部执行完成。
52+ 
Adocs/02_dev_guide/系统任务.md+117-0
@@ -0,0 +1,117 @@
1+# 系统任务
2+ 
3+除了可以下发Kernel执行任务外，Runtime还提供下发Reduce和随机数生成的内置系统任务的功能。（系统任务区别于Kernel任务在于无需用户提供执行代码）。 系统任务可以下发到某条Stream异步执行，同样遵循同一流上任务保序执行的规则。
4+ 
5+通过aclrtReduceAsync接口可以下发执行Reduce操作任务，调用代码示例如下：
6+ 
7+```
8+ aclInit(NULL);
9+ aclrtSetDevice(0);
10+ aclrtStream stream;
11+ aclrtCreateStream(&stream);
12+ // 准备 Host 数据
13+ const int count = 4;
14+ float hostInput[4] = {1.0, 2.0, 3.0, 4.0};
15+ float hostOutput[4] = {0, 0, 0, 0};
16+ size_t size = count * sizeof(float);
17+ // 申请 Device 内存
18+ void *devInput = NULL;
19+ void *devOutput = NULL;
20+ aclrtMalloc(&devInput, size, ACL_MEM_MALLOC_HUGE_FIRST);
21+ aclrtMalloc(&devOutput, size, ACL_MEM_MALLOC_HUGE_FIRST);
22+ // 拷贝数据到 Device
23+ aclrtMemcpy(devInput, size, hostInput, size, ACL_MEMCPY_HOST_TO_DEVICE);
24+ aclrtMemcpy(devOutput, size, hostInput, size, ACL_MEMCPY_HOST_TO_DEVICE);
25+ // 调用 aclrtReduceAsync
26+ aclrtReduceAsync(devOutput,
27+ devInput,
28+ size,
29+ ACL_RT_MEMCPY_SDMA_AUTOMATIC_SUM, // 归约类型
30+ ACL_FLOAT, // 数据类型
31+ stream,
32+ NULL);
33+ // 同步 stream
34+ aclrtSynchronizeStream(stream);
35+ // 拷回结果
36+ aclrtMemcpy(hostOutput, size, devOutput, size, ACL_MEMCPY_DEVICE_TO_HOST);
37+ for (int i = 0; i < count; i++) {
38+ printf("Reduce SUM result[%d] = %f\n", i, hostOutput[i]);
39+ }
40+ /* 预期如下结果
41+ Reduce SUM result[0] = 2.000000
42+ Reduce SUM result[1] = 4.000000
43+ Reduce SUM result[2] = 6.000000
44+ Reduce SUM result[3] = 8.000000
45+ */
46+ // 释放资源
47+ aclrtFree(devInput);
48+ aclrtFree(devOutput);
49+ aclrtDestroyStream(stream);
50+ aclrtResetDeviceForce(0);
51+ aclFinalize();
52+```
53+ 
54+通过aclrtRandomNumAsync执行随机数生成任务，调用代码示例如下：
55+ 
56+```
57+aclError NormalFloatAsync(
58+ float mean, float stddev, uint64_t seed, uint64_t num, void *counterDevAddr, void *devOutput, aclrtStream stream)
59+{
60+ aclrtRandomNumTaskInfo taskInfo = {};
61+ taskInfo.dataType = ACL_FLOAT;
62+ taskInfo.randomNumFuncParaInfo.funcType = ACL_RT_RANDOM_NUM_FUNC_TYPE_NORMAL_DIS;
63+ taskInfo.randomParaAddr = NULL;
64+ taskInfo.randomCounterAddr = counterDevAddr;
65+ taskInfo.randomResultAddr = devOutput;
66+ memcpy(taskInfo.randomNumFuncParaInfo.paramInfo.normalDisInfo.mean.valueOrAddr, &mean, sizeof(float));
67+ taskInfo.randomNumFuncParaInfo.paramInfo.normalDisInfo.mean.size = sizeof(float);
68+ taskInfo.randomNumFuncParaInfo.paramInfo.normalDisInfo.mean.isAddr = 0;
69+ memcpy(taskInfo.randomNumFuncParaInfo.paramInfo.normalDisInfo.stddev.valueOrAddr, &stddev, sizeof(float));
70+ taskInfo.randomNumFuncParaInfo.paramInfo.normalDisInfo.stddev.size = sizeof(float);
71+ taskInfo.randomNumFuncParaInfo.paramInfo.normalDisInfo.stddev.isAddr = 0;
72+ memcpy(taskInfo.randomSeed.valueOrAddr, &seed, sizeof(uint64_t));
73+ taskInfo.randomSeed.size = sizeof(uint64_t);
74+ taskInfo.randomSeed.isAddr = 0;
75+ memcpy(taskInfo.randomNum.valueOrAddr, &num, sizeof(uint64_t));
76+ taskInfo.randomNum.size = sizeof(uint64_t);
77+ taskInfo.randomNum.isAddr = 0;
78+ return aclrtRandomNumAsync(&taskInfo, stream, NULL);
79+}
80+int main()
81+{
82+ aclError ret;
83+ // 初始化 ACL
84+ ret = aclInit(NULL);
85+ ret = aclrtSetDevice(0);
86+ aclrtStream stream;
87+ ret = aclrtCreateStream(&stream);
88+ uint64_t num = 128;
89+ size_t size = num * sizeof(uint64_t); // 申请足够大内存
90+ // 申请 Device 内存
91+ void *devOutput = NULL;
92+ ret = aclrtMalloc(&devOutput, size, ACL_MEM_MALLOC_NORMAL_ONLY);
93+ // 准备 Host 数据
94+ void *hostOutput = malloc(size);
95+ // 申请存放随机数状态 counter 的device内存 （要求 16Byte)
96+ void *counterAddr = NULL;
97+ ret = aclrtMalloc((void **)&counterAddr, 16, ACL_MEM_MALLOC_NORMAL_ONLY);
98+
99+ float mean = 3.0;
100+ float stddev = 2.0;
101+ ret =NormalFloatAsync(mean, stddev, 0, num, counterAddr, devOutput, stream);
102+ // 同步 stream
103+ aclrtSynchronizeStream(stream);
104+ // 拷回结果
105+ aclrtMemcpy(hostOutput, size, devOutput, size, ACL_MEMCPY_DEVICE_TO_HOST);
106+
107+ // 释放资源
108+ free(hostOutput);
109+ aclrtFree(devOutput);
110+ aclrtFree(counterAddr);
111+ aclrtDestroyStream(stream);
112+ aclrtResetDeviceForce(0);
113+ aclFinalize();
114+ return 0;
115+}
116+```
117+ 
Adocs/02_dev_guide/获取Runtime错误码.md+44-0
@@ -0,0 +1,44 @@
1+# 获取Runtime错误码
2+ 
3+所有Runtime接口都会返回一个错误码。然而，对于异步接口（参见[异步任务执行](异步任务执行.md)），由于接口在Device任务完成前就已经返回，因此无法报告Device上可能发生的异步任务错误。它只能返回在Device任务执行前发生在Host上的错误，例如参数校验失败。如果发生异步任务错误，对应的错误码将在后续某个无关的Runtime接口调用时返回。
4+ 
5+因此，若要在某个异步函数调用后立即检查异步错误，唯一的方法是在调用该函数后立即调用aclrtSynchronizeDevice接口（或使用[显式同步](显式同步.md)中描述的任何其他同步机制）进行同步，并检查aclrtSynchronizeDevice接口返回的错误码。
6+ 
7+Runtime会为每个Host线程维护一个错误变量，该变量初始化为ACL\_RT\_SUCCESS，并在每次发生错误（无论是参数校验错误还是异步错误）时被错误码覆盖。aclrtPeekAtLastError接口会返回这个变量的值。aclrtGetLastError接口也会返回这个变量，但同时会将其重置为ACL\_RT\_SUCCESS。
8+ 
9+```
10+// 指定Device，所有的runtime函数都会返回其错误码，可以校验其执行是否成功
11+aclError error = aclrtSetDevice(0);
12+ 
13+// 创建Stream
14+aclrtStream stream;
15+error = aclrtCreateStream(&stream);
16+ 
17+// 设置遇错即停模式
18+aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE);
19+ 
20+// 在Stream上下发任务，返回码仅表示下发是否成功，通常是host上参数校验错误，无法表示device上的实际执行错误
21+error = aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize, ACL_MEMCPY_HOST_TO_DEVICE, stream);
22+error = myKernel<<<8, nullptr, stream>>>();
23+error = aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize, ACL_MEMCPY_DEVICE_TO_HOST, stream);
24+ 
25+// 阻塞应用程序运行，直到正在运算中的Device完成运算，并获取当前异步任务的错误码
26+error = aclrtSynchronizeDevice();
27+ 
28+// 可以获取到当前线程最近的发生的错误
29+error = aclrtPeekAtLastError(ACL_RT_THREAD_LEVEL);
30+ 
31+// 获取当前线程最近的错误并将状态重置为ACL_RT_SUCCESS
32+error = aclrtGetLastError(ACL_RT_THREAD_LEVEL);
33+ 
34+// 获取ErrorMsg，输出到日志中
35+char *errMsg = aclGetRecentErrMsg();
36+......
37+ 
38+// 资源销毁
39+error = aclrtDestroyStream(stream);
40+error = aclrtResetDevice(0);
41+```
42+ 
43+**注意：**在遇错继续模式下（具体请参见[配置任务遇错即停](配置任务遇错即停.md)），如果一条Stream上的任务执行出现异常，该Stream上的其他未执行任务仍可继续执行，同时也不会阻止向该Stream及同处于同一Context下的其他Stream下发新任务。此时，aclrtPeekAtLastError和aclrtGetLastError返回的可能不是首次错误的信息。
44+ 
Adocs/02_dev_guide/虚拟内存管理.md+64-0
@@ -0,0 +1,64 @@
1+# 虚拟内存管理
2+ 
3+Runtime提供了一套虚拟内存管理的API接口，可供开发者更加精细化的管理其内存，例如：可以提前申请好大片的连续的虚拟内存，以简化后续的管理使用，同时可以按需在使用过程中申请物理内存并与虚拟内存做地址映射，即可以精细化管理，也可以最大化利用物理内存；此外还可以在进程之间利用虚拟内存管理中的handle来实现物理内存共享。
4+ 
5+以下是一个虚拟内存的基本使用示例，申请好物理内存以及虚拟内存后，通过aclrtMapMem接口进行映射，即可通过虚拟地址访问对应的数据：
6+ 
7+```
8+void SampleTest()
9+{
10+ // 1. 查询内存申请粒度，并对size做对齐处理
11+ // 设置要申请的物理内存属性
12+ aclrtPhysicalMemProp prop = {};
13+ prop.handleType = ACL_MEM_HANDLE_TYPE_NONE; // handle类型，当前仅支持 ACL_MEM_HANDLE_TYPE_NONE
14+ prop.allocationType = ACL_MEM_ALLOCATION_TYPE_PINNED; // 内存分配类型，当前仅支持分配锁页内存
15+ prop.memAttr = ACL_HBM_MEM_NORMAL; // 内存属性，这里申请HBM内存
16+ prop.location.type = ACL_MEM_LOCATION_TYPE_DEVICE; // 内存所在位置，这里申请Device内存，且deviceId为0
17+ prop.location.id = 0;
18+ // 查询申请最小内存粒度
19+ size_t granularity = 0UL;
20+ aclrtMemGetAllocationGranularity(&prop, ACL_RT_MEM_ALLOC_GRANULARITY_MINIMUM, &granularity);
21+ // 根据查询到的内存申请粒度做内存对齐，以便节约内存
22+ const size_t dataSize = 1024 * sizeof(float);
23+ size_t alignedSize = ((dataSize + granularity - 1U) / granularity) * granularity;
24+ 
25+ // 2. 申请物理内存
26+ // 获取存放物理内存信息的handle，用于后续申请虚拟内存使用，需要注意的是handle并不能直接当虚拟地址访问
27+ aclrtDrvMemHandle handle = nullptr;
28+ aclrtMallocPhysical(&handle, alignedSize, &prop, 0); // 最后一个flag参数当前默认填0
29+ 
30+ // 3. 申请预留虚拟内存
31+ void *virPtr;
32+ // virPtr即为申请预留的虚拟地址指针，alignedSize为申请预留的内存大小
33+ aclrtReserveMemAddress(&virPtr, alignedSize, 0, nullptr, 0);
34+ 
35+ // 4. 映射虚拟内存到物理内存上
36+ // 将申请物理内存时获取的handle与申请预留的虚拟地址空间关联起来
37+ // 需要注意的是，映射的内存大小alignedSize要与申请物理内存时设置的size大小一致
38+ aclrtMapMem(virPtr, alignedSize, 0, handle, 0);
39+ 
40+ // 5. 设置虚拟内存访问权限
41+ aclrtMemAccessDesc accessDesc = {};
42+ accessDesc.flags = ACL_RT_MEM_ACCESS_FLAGS_READWRITE;
43+ accessDesc.location.type = ACL_MEM_LOCATION_TYPE_DEVICE;
44+ accessDesc.location.id = 0;
45+ aclrtMemSetAccess(virPtr, alignedSize, &accessDesc, 1);
46+ 
47+ // 6. 内存访问
48+ // 这里以Device->Host的内存拷贝为例，展示正常访问映射后的内存地址
49+ int *hostPtr;
50+ aclrtMallocHost(reinterpret_cast<void**>(&hostPtr), alignedSize);
51+ aclrtMemcpy(hostPtr, alignedSize, virPtr, alignedSize, ACL_MEMCPY_DEVICE_TO_HOST);
52+ aclrtFreeHost(hostPtr);
53+ 
54+ // 7. 取消虚拟内存与物理内存的映射
55+ aclrtUnmapMem(virPtr);
56+ 
57+ // 8. 释放虚拟内存
58+ aclrtReleaseMemAddress(virPtr);
59+ 
60+ // 9. 释放物理内存
61+ aclrtFreePhysical(handle);
62+}
63+```
64+ 
Adocs/02_dev_guide/记录Event时间戳.md+30-0
@@ -0,0 +1,30 @@
1+# 记录Event时间戳
2+ 
3+在[Event的创建与销毁](Event的创建与销毁.md)章节中创建的Event可用于统计Stream上计算任务的耗时，代码示例如下。该示例仅用于说明Event使用方法，不可以直接拷贝编译运行。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/blob/master/example/event/1_event_timestamp)。
4+ 
5+```
6+uint64_t time = 0;
7+float useTime = 0;
8+ 
9+// 创建Stream
10+aclrtStream stream;
11+aclrtCreateStream(&stream);
12+ 
13+aclrtEvent startEvent;
14+aclrtEvent endEvent;
15+// 创建Event，接口传入ACL_EVENT_TIME_LINE参数，表示创建的Event用于记录
16+aclrtCreateEventExWithFlag(&startEvent, ACL_EVENT_TIME_LINE);
17+aclrtCreateEventExWithFlag(&endEvent, ACL_EVENT_TIME_LINE);
18+ 
19+// 插入startEvent
20+aclrtRecordEvent(startEvent, stream);
21+// 在Stream中下发计算任务
22+kernel<<< grid, block, 0, stream>>>(...);
23+// 插入endEvent
24+aclrtRecordEvent(endEvent, stream);
25+aclrtSynchronizeStream(stream);
26+ 
27+// 获取时间戳并计算耗时
28+aclrtEventElapsedTime(&useTime, startEvent, endEvent);
29+```
30+ 
Adocs/02_dev_guide/跨Device的数据交互.md+35-0
@@ -0,0 +1,35 @@
1+# 跨Device的数据交互
2+ 
3+本节中的“跨Device的数据交互”是指一个进程内、根据硬件组网（例如处于PCIe或者HCCS互联的组网拓扑下）、Device之间能够访问彼此的内存。可以使用aclrtDeviceCanAccessPeer接口查询两个Device之间是否支持数据交互，若支持，再根据访问方向，分别调用aclrtDeviceEnablePeerAccess接口开启一个Device到另一个Device的数据交互功能，例如，调用一次aclrtDeviceEnablePeerAccess接口开启Device 0到Device 1的数据交互，再调用一次aclrtDeviceEnablePeerAccess接口开启Device 1到Device 0的数据交互。若需关闭Device之间的数据交互，可调用aclrtDeviceDisablePeerAccess接口。对于两个进程之间的通信请参见[进程间通信](进程间通信.md)。
4+ 
5+以下是跨Device内存复制的代码示例，不可以直接拷贝编译运行，仅供参考。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/tree/master/example/device/2_device_P2P)。
6+ 
7+```
8+aclInit(NULL); // 初始化
9+int32_t canAccessPeer = 0;
10+aclrtDeviceCanAccessPeer(&canAccessPeer, 0, 1); // 查询Device 0和Device 1之间是否支持数据交互
11+if (canAccessPeer == 1) {
12+ aclrtSetDevice(0); // Device 0下的操作
13+ uint32_t reserveFlag = 0U;
14+ aclrtDeviceEnablePeerAccess(1, reserveFlag); // 开启当前Device(Device 0)到指定Device(Device 1）的数据交互
15+ void *dev0Mem = nullptr;
16+ aclrtMalloc(&dev0Mem, 10, ACL_MEM_MALLOC_HUGE_FIRST_P2P);
17+ aclrtSetDevice(1); // Device 1下的操作
18+ aclrtDeviceEnablePeerAccess(0, reserveFlag); // 开启当前Device(Device 1)到指定Device(Device 0）的数据交互
19+ void *dev1Mem = nullptr;
20+ aclrtMalloc(&dev1Mem, 10, ACL_MEM_MALLOC_HUGE_FIRST_P2P);
21+ aclrtMemcpy(dev1Mem, 10, dev0Mem, 10, ACL_MEMCPY_DEVICE_TO_DEVICE); // 将Device 0上的内存数据复制到Device 1上
22+ aclrtDeviceDisablePeerAccess(0); // 关闭当前Device（Device 1）到指定Device（Device 0）的数据交互
23+ aclrtFree(dev1Mem);
24+ aclrtResetDevice(1); // 释放Device 1的资源
25+ aclrtSetDevice(0); // 切换到Device 0上进行操作
26+ aclrtDeviceDisablePeerAccess(1); // 关闭当前Device（Device 0）到指定Device（Device 1）的数据交互
27+ aclrtFree(dev0Mem);
28+ aclrtResetDeviceForce(0); // Device 0下的操作，调用aclrtResetDevice接口释放Device 0的资源
29+ printf("P2P copy success\n");
30+} else {
31+ printf("current device doesn't support p2p feature\n");
32+}
33+aclFinalize(); // 去初始化
34+```
35+ 
Adocs/02_dev_guide/跨流捕获.md+171-0
@@ -0,0 +1,171 @@
1+# 跨流捕获
2+ 
3+**须知：**本功能为试验特性，后续版本可能会存在变更，不支持应用于商用产品中。
4+ 
5+在捕获Stream上的任务时，aclmdlRICaptureBegin和aclmdlRICaptureEnd接口中指定的Stream只能是同一个Stream（我们将其称为“主流”），若要实现跨Stream的捕获任务，可调用aclrtRecordEvent接口在主流上下发Event Record任务、在其他Stream上调用aclrtStreamWaitEvent接口下发Event Wait任务，以建立主流与其他Stream的关联关系，从而将主流以及其他Stream上的任务捕获到同一个模型中。这时，该Event也会进入捕获状态，如果还有其他Stream等待该Event，那么，相应的Stream也会进入捕获状态。如下图所示，Stream2需等待主流上的task1任务完成，Stream3需等待Stream2上的task2任务完成，这种情况下，Stream2直接依赖主流，而Stream3相当于间接依赖主流，因此，Stream2、Stream3均会被纳入捕获状态，Stream2上的task2任务、Stream3上的task3任务也会被捕获到模型中。
6+ 
7+通过Event加入捕获状态的Stream，最终还需要直接或间接地再通过Event返回到主流，否则会在结束捕获时触发报错。如下图所示，可调用aclrtRecordEvent接口在Stream2、Stream3上下发Event Record任务、在主流上调用aclrtStreamWaitEvent接口下发Event Wait任务，以实现Stream2、Stream3返回主流。另外，对于像Stream3这种间接依赖主流的情况，也可以在Stream3上下发Event Record任务、在Stream2上下发Event Wait任务，先将Stream3返回Stream2，然后再在Stream2上下发Event Record任务、在主流上下发Event Wait任务，最终返回到主流。返回主流之后，结束捕获前，不能再在Stream2、Stream3上下发task（例如下图中的task5），否则在结束捕获时会因为校验到有未被关联的task而触发报错。
8+ 
9+跨Stream的任务捕获流程如下图所示：
10+ 
11+![](figures/ACL_Graph跨流捕获.png)
12+ 
13+以下示例用两个stream为例演示跨流捕获，其中stream1是主流。
14+ 
15+```
16+#include <stdio.h>
17+#include <vector>
18+#include "acl/acl.h"
19+#include "aclnnop/aclnn_add.h"
20+ 
21+#define ACL_LOG(fmt, args...) fprintf(stdout, "[INFO] " fmt "\n", ##args)
22+ 
23+int64_t GetShapeSize(const std::vector<int64_t> &shape)
24+{
25+ int64_t shape_size = 1;
26+ for (auto i : shape) {
27+ shape_size *= i;
28+ }
29+ return shape_size;
30+}
31+ 
32+int CreateAclTensor(const std::vector<int64_t> &shape, void **deviceAddr,
33+ aclDataType dataType, aclTensor **tensor)
34+{
35+ auto size = GetShapeSize(shape) * sizeof(float);
36+ // 申请Device侧内存
37+ auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST);
38+ // 计算连续tensor的stride
39+ std::vector<int64_t> strides(shape.size(), 1);
40+ for (int64_t i = shape.size() - 2; i >= 0; i--) {
41+ strides[i] = shape[i + 1] * strides[i + 1];
42+ }
43+ // 调用aclCreateTensor接口创建aclTensor
44+ *tensor = aclCreateTensor(shape.data(),
45+ shape.size(),
46+ dataType,
47+ strides.data(),
48+ 0,
49+ aclFormat::ACL_FORMAT_ND,
50+ shape.data(),
51+ shape.size(),
52+ *deviceAddr);
53+ return 0;
54+}
55+ 
56+int main()
57+{
58+ int devID = 0;
59+ void *self_d = nullptr;
60+ void *other_d = nullptr;
61+ void *out_d = nullptr;
62+ aclTensor *self = nullptr;
63+ aclTensor *other = nullptr;
64+ aclScalar *alpha = nullptr;
65+ aclTensor *out = nullptr;
66+ /* aclnnAdd: out = self + other * alpha */
67+ float *self_h = nullptr;
68+ float *other_h = nullptr;
69+ std::vector<int64_t> shape = {4, 2};
70+ float alphaValue = 1.1f;
71+ uint64_t workspaceSize = 0;
72+ aclOpExecutor *executor;
73+ auto size = GetShapeSize(shape);
74+
75+ // 初始化
76+ aclInit(NULL);
77+ // 指定计算设备
78+ aclrtSetDevice(devID);
79+ 
80+ // 准备aclnnAdd算子的输入、输出参数
81+ CreateAclTensor(shape, &self_d, aclDataType::ACL_FLOAT, &self);
82+ CreateAclTensor(shape, &other_d, aclDataType::ACL_FLOAT, &other);
83+ alpha = aclCreateScalar(&alphaValue, aclDataType::ACL_FLOAT);
84+ CreateAclTensor(shape, &out_d, aclDataType::ACL_FLOAT, &out);
85+ 
86+ // 获取算子计算所需的workspace大小以及包含了算子计算流程的执行器
87+ aclnnAddGetWorkspaceSize(self, other, alpha, out, &workspaceSize, &executor);
88+ void *workspaceAddr = nullptr;
89+ if (workspaceSize > 0) {
90+ aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
91+ }
92+ 
93+ // 使用aclrtMallocHost申请锁页内存
94+ aclrtMallocHost((void **)&self_h, size * sizeof(float));
95+ aclrtMallocHost((void **)&other_h, size * sizeof(float));
96+ for (int i = 0; i < 8; i++) {
97+ self_h[i] = static_cast<float>(0);
98+ other_h[i] = static_cast<float>(1);
99+ }
100+ 
101+ aclmdlRI modelRI;
102+ aclrtStream stream1, stream2;
103+ aclrtEvent event1, event2;
104+ aclrtCreateStream(&stream1);
105+ aclrtCreateStream(&stream2);
106+ aclrtCreateEvent(&event1);
107+ aclrtCreateEvent(&event2);
108+ 
109+ // ========开始捕获任务========
110+ aclmdlRICaptureBegin(stream1, ACL_MODEL_RI_CAPTURE_MODE_GLOBAL);
111+ // 异步拷贝，将算子self输入的数据从Host侧传到Device侧
112+ aclrtMemcpyAsync(self_d, size * sizeof(float), self_h, size * sizeof(float), ACL_MEMCPY_HOST_TO_DEVICE, stream1);
113+ // 切换捕获模式为RELAXED，允许调用aclrtMemcpy函数
114+ aclmdlRICaptureMode mode = ACL_MODEL_RI_CAPTURE_MODE_RELAXED;
115+ aclmdlRICaptureThreadExchangeMode(&mode);
116+ // 同步拷贝，将算子other输入的数据从Host侧传到Device侧，仅执行一次
117+ aclrtMemcpy(other_d, size * sizeof(float), other_h, size * sizeof(float), ACL_MEMCPY_HOST_TO_DEVICE);
118+ // 将捕获模式切换回GLOBAL
119+ aclmdlRICaptureThreadExchangeMode(&mode);
120+ // 通过event1，将stream2加入捕获状态
121+ aclrtRecordEvent(event1, stream1);
122+ aclrtStreamWaitEvent(stream2, event1);
123+ // 执行aclnnAdd算子
124+ aclnnAdd(workspaceAddr, workspaceSize, executor, stream2);
125+ // stream2上的任务执行完成后，通过event2，让stream2返回主流stream1
126+ aclrtRecordEvent(event2, stream2);
127+ aclrtStreamWaitEvent(stream1, event2);
128+ // 异步拷贝，将算子输出数据从Device侧传回Host侧
129+ aclrtMemcpyAsync(self_h, size * sizeof(float), out_d, size * sizeof(float), ACL_MEMCPY_DEVICE_TO_HOST, stream1);
130+ // ========结束捕获任务========
131+ aclmdlRICaptureEnd(stream1, &modelRI);
132+ 
133+ // 多次执行模型
134+ for (int i = 0; i < 8; i++) {
135+ aclmdlRIExecuteAsync(modelRI, stream1);
136+ aclrtSynchronizeStream(stream1);
137+ // 打印每一次的算子输出数据
138+ ACL_LOG("%f %f %f %f %f %f %f %f\n",
139+ self_h[0],
140+ self_h[1],
141+ self_h[2],
142+ self_h[3],
143+ self_h[4],
144+ self_h[5],
145+ self_h[6],
146+ self_h[7]);
147+ }
148+ 
149+ // 释放资源
150+ aclmdlRIDestroy(modelRI);
151+ aclrtDestroyStream(stream1);
152+ aclrtDestroyStream(stream2);
153+ aclrtDestroyEvent(event1);
154+ aclrtDestroyEvent(event2);
155+ aclDestroyTensor(self);
156+ aclDestroyTensor(other);
157+ aclDestroyTensor(out);
158+ aclDestroyScalar(alpha);
159+ aclrtFree(self_d);
160+ aclrtFree(other_d);
161+ aclrtFree(out_d);
162+ if (workspaceAddr != nullptr) {
163+ aclrtFree(workspaceAddr);
164+ }
165+ // 释放计算设备的资源
166+ aclrtResetDevice(devID);
167+ // 去初始化
168+ aclFinalize();
169+}
170+```
171+ 
Adocs/02_dev_guide/运行时核资源控制.md+76-0
@@ -0,0 +1,76 @@
1+# 运行时核资源控制
2+ 
3+为了提高Device核资源的使用率以及隔离性，Runtime支持控制Device的核资源。当前支持配置Device粒度、Stream粒度的核资源限制，核资源包括AI Core或Cube Core数量、Vector Core数量。
4+ 
5+是否由用户显式指定numBlocks（用于指定算子的核函数将会在几个核上执行），调用运行时核资源控制的接口会有所不同，如下所示：
6+ 
7+- **无numBlocks场景**（例如调用aclnn算子接口）：Stream粒度的核资源限制需要调用aclrtUseStreamResInCurrentThread接口绑定到当前线程使用，并调用aclrtGetResInCurrentThread接口获取当前线程可使用的核资源。
8+ 
9+ 通过aclrtGetResInCurrentThread接口获取核资源限制的优先级为：Stream粒度的核资源限制 \> Device粒度的核资源限制 \> AI处理器硬件的默认核资源限制。例如，Device总共包含32个Vector Core，Device粒度限制使用16个Vector Core，而Stream粒度的核资源限制可以为20个Vector Core，则aclnn算子执行时以20个Vector Core运行。
10+ 
11+- **需numBlocks场景**（例如LaunchKernel方式执行算子）：用户可调用aclrtGetDeviceResLimit接口、aclrtGetStreamResLimit接口获取不同粒度的核资源限制后再配置numBlocks。
12+ 
13+## Device粒度的核资源限制
14+ 
15+以下是关键步骤的代码示例，不可以直接拷贝编译运行，仅供参考。完整样例代码，请参见[Link](https://gitcode.com/cann/runtime/tree/master/example/kernel/1_launch_kernel_with_reslimit)。
16+ 
17+```
18+......
19+int32_t deviceId = 0;
20+uint32_t numBlocks = 8;
21+uint32_t coreDim = 0;
22+ 
23+// 指定计算设备
24+aclrtSetDevice(deviceId);
25+ 
26+// 设置核资源限制
27+aclrtSetDeviceResLimit(deviceId, ACL_RT_DEV_RES_CUBE_CORE, numBlocks);
28+ 
29+// 下发算子执行任务
30+// 以aclnnAdd算子为例：
31+// aclnnAddGetWorkspaceSize中会隐式调用aclrtGetResInCurrentThread查询核资源进行tiling
32+// aclnnAdd()
33+......
34+ 
35+// 获取核资源限制
36+aclrtGetDeviceResLimit(deviceId, ACL_RT_DEV_RES_CUBE_CORE, &coreDim);
37+```
38+ 
39+## Stream粒度的核资源限制
40+ 
41+以下是关键步骤的代码示例，不可以直接拷贝编译运行，仅供参考。
42+ 
43+```
44+#include "acl/acl.h"
45+......
46+int32_t deviceId = 0;
47+uint32_t numBlocks = 8;
48+ 
49+// 指定运算的Device
50+aclrtSetDevice(deviceId);
51+ 
52+// 显式创建一个Stream
53+aclrtStream stream;
54+aclrtCreateStream(&stream);
55+ 
56+// 设置Stream粒度的Cube Core、Vector Core的数量
57+aclrtSetStreamResLimit(stream, ACL_RT_DEV_RES_CUBE_CORE, numBlocks);
58+aclrtSetStreamResLimit(stream, ACL_RT_DEV_RES_VECTOR_CORE, numBlocks * 2);
59+ 
60+// 绑定到当前线程
61+aclrtUseStreamResInCurrentThread(stream);
62+ 
63+// 算子的任务中需要调用aclrtGetResInCurrentThread查询当前线程的资源限制，然后指定运行算子的资源数量
64+// 以aclnnAdd算子为例：
65+// aclnnAddGetWorkspaceSize中会隐式调用aclrtGetResInCurrentThread查询核资源进行tiling
66+// aclnnAdd()
67+......
68+ 
69+aclrtUnuseStreamResInCurrentThread(stream);
70+aclrtResetStreamResLimit(stream);
71+ 
72+// 资源销毁
73+aclrtDestroyStream(stream);
74+aclrtResetDevice(deviceId);
75+```
76+ 
Adocs/02_dev_guide/进程间通信.md+278-0
@@ -0,0 +1,278 @@
1+# 进程间通信
2+ 
3+由某个主机线程创建的任意设备内存、Event资源或Notify资源，都可以在同一进程内被该进程中的其他线程直接引用。但这些指针或句柄在进程之外是无效的，因此不能被其他进程的线程直接使用。
4+ 
5+若要在不同进程之间共享设备内存、Event资源或Notify资源，需要应用程序使用Runtime提供的进程间通信相关API以实现如下典型场景：**由一个主进程生成一批输入数据，并将这些数据提供给多个从属进程使用，而无需在每个进程中重新生成或复制数据**。不同资源涉及的IPC（Inter-Process Communication）接口不同，可查看下文中的调用示例。
6+ 
7+需要注意的是，通过aclrtMalloc接口分配设备内存时，出于性能考虑，可能会从更大的底层内存块中切分出来。在这种情况下，IPC接口会检查共享内存是否页表对齐，若未对齐，API将拦截并报错，以防止跨进程多映射内存导致的信息泄露风险。因此，建议使用aclrtMalloc接口根据内存分配规则申请内存。申请不同类型的内存时，其页表大小会有所不同：普通页内存的页表大小为4K，大页内存的页表大小支持2M或1G。
8+ 
9+## 进程间共享内存
10+ 
11+此处以A进程（内存出借方）、B进程（内存借入方）为例，说明两个进程间的内存共享接口调用流程：
12+ 
13+![](figures/进程间共享内存.png)
14+ 
15+以下为A、B进程之间共享内存的示例代码，不可以直接拷贝编译运行，仅供参考。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/tree/master/example/memory/11_ipc_memory_withoutpid)。
16+ 
17+1. 在A进程中分配内存，生成共享key：
18+ 
19+ ```
20+ uint keyLen = 65;
21+ char[keyLen] key;
22+ void *ptrA = nullptr;
23+
24+ aclrtSetDevice(0); // 进程A使用Device 0
25+ aclrtMalloc(&ptrA, size);
26+
27+ aclrtIpcMemGetExportKey(ptrA, size, key, keyLen, ACL_RT_IPC_MEM_EXPORT_FLAG_DISABLE_PID_VALIDATION);
28+
29+ // 跨进程通信交换key（以写文件方式交互)
30+ writeFile("file/ipc_mem", key, keyLen);
31+
32+ // 对共享内存进行读写操作
33+ ......
34+
35+ // 待共享内存使用完成后，借出方关闭IPC共享内存
36+ aclrtIpcMemClose(key);
37+ aclrtFree(ptrA);
38+ aclrtResetDeviceForce(0);
39+ ```
40+ 
41+2. 在B进程中，通过共享key导入共享内存：
42+ 
43+ ```
44+ uint keyLen = 65;
45+ char[keyLen] key;
46+ void *ptrB;
47+
48+ // 跨进程通信交换key（以写文件方式交互)
49+ readFile("file/ipc_mem", key, keyLen);
50+
51+ aclrtSetDevice(1); // 进程B使用Device 1
52+
53+ // A，B进程使用不同的Device，导入共享内存访问要开启两个Device之间的数据交互
54+ aclrtIpcMemImportByKey(&ptrB, key, ACL_RT_IPC_MEM_IMPORT_FLAG_ENABLE_PEER_ACCESS);
55+
56+ // 对ptrB内存读写操作
57+ ......
58+
59+ // 使用完成后，借入方关闭IPC共享内存
60+ aclrtIpcMemClose(key);
61+ aclrtResetDeviceForce(1);
62+ ```
63+ 
64+## 进程间共享Event
65+ 
66+进程之间通过共享Event，可以实现进程间的事件同步。此处以A进程创建Event，共享给B进程为例，说明两个进程间任务同步的示例代码，不可以直接拷贝编译运行，仅供参考。
67+ 
68+1. 在A进程中创建Event，生成共享handle：
69+ 
70+ ```
71+ aclrtEvent event;
72+ aclrtStream stream;
73+ aclrtIpcEventHandle handle;
74+
75+ aclrtSetDevice(0); // 进程A使用Device 0
76+ aclrtCreateEventExWithFlag(&event, ACL_EVENT_IPC); // 创建IPC Event
77+ aclrtCreateStream(&stream); // 创建Stream
78+
79+ // 导出进程间共享handle
80+ aclrtIpcGetEventHandle(event, &handle);
81+
82+ // 跨进程通信交换key（以写文件方式交互)
83+ writeFile("file/ipc_event", handle, ACL_IPC_EVENT_HANDLE_SIZE);
84+
85+ // 下发record任务
86+ aclrtEventRecord(event, stream);
87+
88+ // Event使用完, 销毁共享Event
89+ aclrtDestroyEvent(event);
90+ aclrtResetDeviceForce(0);
91+ ```
92+ 
93+2. 在B进程中，通过共享handle导入共享Event：
94+ 
95+ ```
96+ aclrtStream stream;
97+ aclrtEvent event;
98+ aclrtIpcEventHandle handle;
99+
100+ aclrtCreateStream(&stream); // 创建Stream
101+ // 跨进程获取共享handle（以写文件方式交互)
102+ readFile("file/ipc_event", handle, ACL_IPC_EVENT_HANDLE_SIZE);
103+
104+ aclrtSetDevice(1); //进程B使用Device 1
105+
106+ // 导入handle，返回共享event
107+ // A，B进程使用不同的Device
108+ aclrtIpcOpenEventHandle(handle, &event);
109+
110+ // 下发wait任务
111+ aclrtStreamWaitEvent(stream, event);
112+
113+ // 同步Stream上的任务
114+ aclrtSynchonizeStream(stream);
115+
116+ // Event使用完，销毁共享Event
117+ aclrtDestroyEvent(event);
118+ aclrtResetDeviceForce(1);
119+ ```
120+ 
121+## 进程间共享Notify
122+ 
123+进程之间通过共享Notify，可以实现进程间的通知。此处以A进程创建Notify，共享给B进程为例，说明两个进程间任务同步的示例代码，不可以直接拷贝编译运行，仅供参考。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/tree/master/example/notify/1_ipc_notify_withoutpid)。
124+ 
125+注意：创建端会分配Notify硬件资源，因此只能由创建端硬件进行Wait。为此共享Notify有使用约束，只能在创建端调aclrtNotifyWait进行Wait，不能在共享端Wait。
126+ 
127+1. 在A进程中创建Notify，生成共享key：
128+ 
129+ ```
130+ uint keyLen = 65;
131+ char key[keyLen];
132+ aclrtNotify notify;
133+ aclrtStream stream;
134+
135+ // 进程A使用Device 0
136+ aclrtSetDevice(0);
137+ aclrtCreateStream(&stream);
138+
139+ aclrtNotifyCreate(&notify);
140+ // 导出key（即Notify共享名称）
141+ aclrtNotifyGetExportKey(notify, key, keyLen, ACL_RT_NOTIFY_EXPORT_FLAG_DISABLE_PID_VALIDATION);
142+ // 跨进程通信交换key（以写文件方式交互)
143+ writeFile("file/ipc_notify", key, keyLen);
144+
145+ // 下发wait任务
146+ aclrtNotifyWait(notify, stream);
147+
148+ // Notify使用完, 销毁共享Notify
149+ aclrtNotifyDestroy(notify);
150+ aclrtResetDeviceForce(0);
151+ ```
152+ 
153+2. 在B进程中，通过共享key导入共享Notify：
154+ 
155+ ```
156+ uint keyLen = 65;
157+ char key[keyLen];
158+ aclrtNotify notify;
159+ aclrtStream stream;
160+
161+ // 跨进程通信交换key（以写文件方式交互)
162+ readFile("file/ipc_notify", key, keyLen);
163+
164+ //进程B使用Device 1
165+ aclrtSetDevice(1);
166+ aclrtCreateStream(&stream);
167+
168+ // 导入key，返回共享Notify
169+ // A，B进程使用不同的Device，导入共享Notify要开启两个Device之间的数据交互
170+ aclrtNotifyImportByKey(&notify, key, ACL_RT_NOTIFY_IMPORT_FLAG_ENABLE_PEER_ACCESS);
171+
172+ // 下发record任务
173+ aclrtNotifyRecord (notify, stream);
174+
175+ // Notify使用完，销毁共享Notify
176+ aclrtNotifyDestroy(notify);
177+ aclrtResetDeviceForce(1);
178+ ```
179+ 
180+## 通过VMM接口实现进程间共享内存
181+ 
182+除IPC Mem共享内存外，Runtime还提供了另外一套内存管理和内存共享接口。VMM（Virtual Memory Management）这套接口提供更灵活的功能，支持虚拟地址申请、物理内存申请和跨进程物理内存共享，还支持虚拟地址与物理内存之间的映射操作。
183+ 
184+此处以A、B进程为例，说明一个Device上、两个进程间的物理内存共享的示例代码，不可以直接拷贝编译运行，仅供参考，完整样例代码请参见[Link](https://gitcode.com/cann/runtime/tree/master/example/memory/8_physical_memory_sharing_withoutpid)。
185+ 
186+若A、B进程使用不同的Device，还需配合使用aclrtDeviceEnablePeerAccess接口开启跨Device的数据交互，详细描述请参见[跨Device的数据交互](跨Device的数据交互.md)。
187+ 
188+1. 在A进程中：
189+ 
190+ ```
191+ // 查询内存申请粒度
192+ const size_t dataSize = 1024 * sizeof(float);
193+ aclrtPhysicalMemProp prop = {};
194+ prop.handleType = ACL_MEM_HANDLE_TYPE_NONE;
195+ prop.allocationType = ACL_MEM_ALLOCATION_TYPE_PINNED;
196+ prop.location.type = ACL_MEM_LOCATION_TYPE_DEVICE;
197+ prop.location.id = 0;
198+ prop.memAttr = ACL_HBM_MEM_NORMAL;
199+ size_t granularity = 0UL;
200+ aclrtMemGetAllocationGranularity(&prop, ACL_RT_MEM_ALLOC_GRANULARITY_MINIMUM, &granularity);
201+
202+ // 基于内存申请粒度申请物理内存
203+ size_t alignedSize = ((dataSize + granularity - 1U) / granularity) * granularity;
204+ aclrtDrvMemHandle handle = nullptr;
205+ aclrtMallocPhysical(&handle, alignedSize, &prop, 0);
206+
207+ // 预留虚拟内存
208+ void *virPtr;
209+ aclrtReserveMemAddress(&virPtr, alignedSize, 0, nullptr, 0);
210+
211+ // 将虚拟内存映射到物理内存
212+ aclrtMapMem(virPtr, alignedSize, 0, handle, 0);
213+ aclrtMemAccessDesc desc = {};
214+ desc.flags = ACL_RT_MEM_ACCESS_FLAGS_READWRITE;
215+ desc.location.id = 0;
216+ desc.location.type = ACL_MEM_LOCATION_TYPE_DEVICE;
217+ aclrtMemSetAccess(virPtr, alignedSize, &desc, 1);
218+
219+ // 使用virPtr进行复制、读、写等操作
220+ ......
221+
222+ // 导入共享handle
223+ uint64_t shareableHandle = 0ULL;
224+ aclrtMemExportToShareableHandle(handle, ACL_MEM_HANDLE_TYPE_NONE, ACL_RT_VMM_EXPORT_FLAG_DISABLE_PID_VALIDATION , &shareableHandle);
225+
226+ // 将共享handle传递给进程B
227+ writeFile("file/vmm_mem", shareableHandle, sizeof(shareableHandle));
228+
229+ // 取消虚拟内存与物理内存之间的映射关系
230+ aclrtUnmapMem(virPtr);
231+ // 释放虚拟内存和物理内存
232+ aclrtReleaseMemAddress(virPtr);
233+ aclrtFreePhysical(handle);
234+ ```
235+ 
236+2. 在B进程中：
237+ 
238+ ```
239+ uint64_t shareableHandle = 0ULL;
240+ // 从文件中获取共享handle
241+ readFile("file/vmm_mem", &shareableHandle, sizeof(shareableHandle));
242+
243+ aclrtDrvMemHandle handle = nullptr;
244+ int32_t deviceId=0;
245+ aclrtMemImportFromShareableHandle(shareableHandle, deviceId, &handle);
246+ // 查询内存申请粒度
247+ const size_t data_size = 1024 * sizeof(float);
248+ aclrtPhysicalMemProp prop = {};
249+ prop.handleType = ACL_MEM_HANDLE_TYPE_NONE;
250+ prop.allocationType = ACL_MEM_ALLOCATION_TYPE_PINNED;
251+ prop.location.type = ACL_MEM_LOCATION_TYPE_DEVICE;
252+ prop.location.id = 0;
253+ prop.memAttr = ACL_HBM_MEM_NORMAL;
254+ size_t granularity = 0UL;
255+ aclrtMemGetAllocationGranularity(&prop, ACL_RT_MEM_ALLOC_GRANULARITY_MINIMUM, &granularity);
256+ size_t alignedSize = ((dataSize + granularity - 1U) / granularity) * granularity;
257+
258+ // 基于内存申请粒度预留虚拟内存
259+ void *virPtr = nullptr;
260+ aclrtReserveMemAddress(&virPtr, alignedSize, 0, nullptr, 0);
261+ // 将虚拟内存映射到物理内存
262+ aclrtMapMem(virPtr, alignedSize, 0, handle, 0);
263+ aclrtMemAccessDesc desc = {};
264+ desc.flags = ACL_RT_MEM_ACCESS_FLAGS_READWRITE;
265+ desc.location.id = 0;
266+ desc.location.type = ACL_MEM_LOCATION_TYPE_DEVICE;
267+ aclrtMemSetAccess(virPtr,alignedSize, &desc, 1);
268+
269+ // 使用virPtr进行复制、读、写等操作
270+ ......
271+
272+ // 取消虚拟内存与物理内存之间的映射关系
273+ aclrtUnmapMem(virPtr);
274+ // 释放虚拟内存和物理内存
275+ aclrtReleaseMemAddress(virPtr);
276+ aclrtFreePhysical(handle);
277+ ```
278+ 
Adocs/02_dev_guide/配置AI-Core栈空间大小.md+16-0
@@ -0,0 +1,16 @@
1+# 配置AI Core栈空间大小
2+ 
3+在CANN中，每个AI Core都拥有自己的私有堆栈，用于存储算子的局部变量和函数调用信息。应用进程在执行aclrtSetDevice时，CANN会为每个AI Core分配大小为aicore\_stack\_size的Device内存作为栈空间，总计为aicore\_stack\_size \* aicore\_num。默认情况下，aicore\_stack\_size = 32KB。
4+ 
5+然而，某些算子在执行时所需的栈空间超过默认的32KB，比如以-O0方式编译用于调试的AI Core算子，其栈空间需求可能会超过32KB。如果不调整系统默认的栈空间大小，可能会导致算子执行异常（如栈溢出）。因此，用户需要在进程启动时，通过修改aclInit初始化接口的json文件来配置合适的AI Core栈空间大小。**注意：**栈空间占用的是Device内存，请按需设置。
6+ 
7+json文件的配置示例如下，在aicore\_stack\_size参数处设置栈空间大小，单位为字节，详细使用说明请参见aclInit接口。
8+ 
9+```
10+{
11+ "StackSize":{
12+ "aicore_stack_size":32768
13+ }
14+}
15+```
16+ 
Adocs/02_dev_guide/配置Stream优先级.md+23-0
@@ -0,0 +1,23 @@
1+# 配置Stream优先级
2+ 
3+在运行时，Device上的调度器会依据各个Stream的优先级来决定任务的执行顺序。高优先级Stream中待执行的任务将优先于低优先级Stream中的任务得到调度，但不会抢占已处于运行状态的低优先级任务。Device在执行过程中不会动态重新评估任务队列，因此提升Stream的优先级不会中断正在执行的任务。
4+ 
5+Stream的优先级主要用于影响任务的调度顺序，而非强制规定严格的执行序列。用户可以通过调整Stream的优先级来引导任务的执行顺序，但无法以此强制保证任务间的绝对顺序。
6+ 
7+调用aclrtCreateStreamWithConfig接口创建Stream时可指定Stream的优先级。允许设置的优先级范围，可以通过aclrtDeviceGetStreamPriorityRange接口获取最小优先级、最大优先级。
8+ 
9+以下为示例代码，不可以直接拷贝编译运行，仅供参考。
10+ 
11+```
12+// 查询当前设备支持的Stream最小、最大优先级
13+aclrtDeviceGetStreamPriorityRange(&leastPriority, &greatestPriority);
14+ 
15+// 创建具有最高和最低优先级的Stream
16+aclrtStream stream_high;
17+aclrtStream stream_low;
18+aclrtCreateStreamWithConfig(&stream_high, greatestPriority, ACL_STREAM_FAST_LAUNCH);
19+aclrtCreateStreamWithConfig(&stream_low, leastPriority, ACL_STREAM_FAST_LAUNCH);
20+```
21+ 
22+Stream的优先级在Device范围内生效，而不是在Context范围内生效。
23+ 
Adocs/02_dev_guide/配置任务遇错即停.md+28-0
@@ -0,0 +1,28 @@
1+# 配置任务遇错即停
2+ 
3+CANN支持遇错即停模式（ACL\_STOP\_ON\_FAILURE）和遇错继续模式（ACL\_CONTINUE\_ON\_FAILURE），以支持不同应用对任务执行失败的差异化控制。默认为遇错继续模式。
4+ 
5+当Stream上的任务执行失败时，如果配置了遇错即停模式（ACL\_STOP\_ON\_FAILURE），则会停止执行该Context中所有Stream上的任务；如果配置了遇错继续模式（ACL\_CONTINUE\_ON\_FAILURE），则会继续执行Stream上的后续任务。
6+ 
7+调用aclrtSetStreamFailureMode接口指定调度模式的示例代码如下，不可以直接拷贝编译运行，仅供参考：
8+ 
9+```
10+aclrtStream stream;
11+aclrtCreateStream(&stream);
12+ 
13+// 设置遇错即停模式
14+aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE);
15+......
16+```
17+ 
18+也可以调用aclrtSetStreamAttribute接口指定调度模式的示例代码如下，不可以直接拷贝编译运行，仅供参考：
19+ 
20+```
21+aclrtStream stream;
22+aclrtCreateStream(&stream);
23+ 
24+// 设置遇错继续模式
25+aclrtSetStreamAttribute(stream, ACL_STREAM_ATTR_FAILURE_MODE, ACL_CONTINUE_ON_FAILURE);
26+......
27+```
28+ 
Adocs/02_dev_guide/默认Stream.md+26-0
@@ -0,0 +1,26 @@
1+# 默认Stream
2+ 
3+在调用aclrtSetDevice接口或aclrtCreateContext接口时，Runtime会自动创建一个默认Stream。每个Context拥有一个默认Stream。如果不同的Host线程使用相同的Context，它们将共享同一个默认Stream。
4+ 
5+对于需要传入Stream参数的API（如aclrtMemcpyAsync），如果使用默认Stream作为入参，则直接传入nullptr。对于没有Stream入参的API（如aclrtMemcpy），则不使用默认Stream。
6+ 
7+默认Stream不能显式调用aclrtDestroyStream接口销毁。在调用aclrtResetDevice或aclrtResetDeviceForce接口释放资源时，默认Stream会被自动销毁。
8+ 
9+以下是在默认Stream上下发计算任务的代码示例，不可以直接拷贝编译运行，仅供参考。
10+ 
11+```
12+// 指定Device（接口内部自动创建默认Stream）
13+aclrtSetDevice(0);
14+ 
15+// 在默认stream上下发Host->Device复制任务、MyKernel任务、和Device->Host复制任务
16+aclrtMemcpyAsync(devPtrIn, size, hostPtr, hostSize, ACL_MEMCPY_HOST_TO_DEVICE, nullptr);
17+myKernel<<<8, nullptr, nullptr>>>(devPtrIn, devPtrOut, size);
18+aclrtMemcpyAsync(hostPtr, hostSize, devPtrOut, size, ACL_MEMCPY_DEVICE_TO_HOST, nullptr);
19+ 
20+// 同步默认stream
21+aclrtStreamSynchronize(nullptr);
22+ 
23+// 复位Device（接口内部自动销毁默认Stream）
24+aclrtResetDevice(0);
25+```
26+ 
Rdocs/api_docs/CntNotify管理.md→docs/03_api_ref/CntNotify管理.md+0-0
文件重命名但无更改。
Rdocs/api_docs/Context管理.md→docs/03_api_ref/Context管理.md+0-0
文件重命名但无更改。
Rdocs/api_docs/Device管理.md→docs/03_api_ref/Device管理.md+0-0
文件重命名但无更改。
Rdocs/api_docs/Dump配置.md→docs/03_api_ref/Dump配置.md+0-0
文件重命名但无更改。
Rdocs/api_docs/Event管理.md→docs/03_api_ref/Event管理.md+0-0
文件重命名但无更改。
Rdocs/api_docs/Kernel加载与执行.md→docs/03_api_ref/Kernel加载与执行.md+0-0
文件重命名但无更改。
Rdocs/api_docs/Label管理.md→docs/03_api_ref/Label管理.md+0-0
文件重命名但无更改。
Rdocs/api_docs/Notify管理.md→docs/03_api_ref/Notify管理.md+0-0
文件重命名但无更改。
Rdocs/api_docs/Profiling数据采集.md→docs/03_api_ref/Profiling数据采集.md+0-0
文件重命名但无更改。
Rdocs/api_docs/Profiling数据采集接口.md→docs/03_api_ref/Profiling数据采集接口.md+0-0
文件重命名但无更改。
Rdocs/api_docs/PyTorch场景标记迭代时间.md→docs/03_api_ref/PyTorch场景标记迭代时间.md+0-0
文件重命名但无更改。
Rdocs/api_docs/SOMA原理说明.md→docs/03_api_ref/SOMA原理说明.md+199-199
@@ -1,199 +1,199 @@
1-# Stream有序内存分配1+# Stream有序内存分配
2- 2+ 
3-## 背景3+## 背景
4-aclrtMalloc和aclrtFree是用于同步内存分配和管理的接口。以下示例代码展示了一个常见的内存使用场景：使用aclrtMalloc申请内存，通过异步拷贝的方式将内存数据拷贝到Device侧以备算子使用，算子执行完成后，通过Stream同步来确认该内存使用完毕，然后使用aclrtFree释放内存。在异步操作较少的情况下，这样的方式是可接受的。然而，在有大量算子下发和对申请内存的高频异步操作时，这种方式存在以下不足：4+aclrtMalloc和aclrtFree是用于同步内存分配和管理的接口。以下示例代码展示了一个常见的内存使用场景：使用aclrtMalloc申请内存，通过异步拷贝的方式将内存数据拷贝到Device侧以备算子使用，算子执行完成后，通过Stream同步来确认该内存使用完毕，然后使用aclrtFree释放内存。在异步操作较少的情况下，这样的方式是可接受的。然而，在有大量算子下发和对申请内存的高频异步操作时，这种方式存在以下不足：
5-1. 同步瓶颈：在任务下发过程中，如果需要进行内存分配或释放的调整，容易产生同步瓶颈，影响整体效率。5+1. 同步瓶颈：在任务下发过程中，如果需要进行内存分配或释放的调整，容易产生同步瓶颈，影响整体效率。
6-2. 累积延迟：内存申请与释放本身耗时，频繁操作会累积不必要的延迟，进一步降低性能。6+2. 累积延迟：内存申请与释放本身耗时，频繁操作会累积不必要的延迟，进一步降低性能。
7- 7+ 
8- 8+ 
9-```cpp9+```cpp
10-#include "acl/acl_rt.h"10+#include "acl/acl_rt.h"
11-#include "acl/acl.h"11+#include "acl/acl.h"
12- 12+ 
13-int main() {13+int main() {
14- typedef struct {14+ typedef struct {
15- ......15+ ......
16- } ArgsInfo;16+ } ArgsInfo;
17- 17+ 
18- void *ptr0 = nullptr;18+ void *ptr0 = nullptr;
19- aclrtStream stream1;19+ aclrtStream stream1;
20- 20+ 
21- // 申请内存21+ // 申请内存
22- aclrtMalloc(&ptr0, sizeof(ArgsInfo), ACL_MEM_MALLOC_HUGE_FIRST);22+ aclrtMalloc(&ptr0, sizeof(ArgsInfo), ACL_MEM_MALLOC_HUGE_FIRST);
23- ......23+ ......
24- 24+ 
25- // 配置任务下发25+ // 配置任务下发
26- ArgsInfo usrArgs;26+ ArgsInfo usrArgs;
27- // 拷贝信息到device侧申请内存27+ // 拷贝信息到device侧申请内存
28- error = aclrtMemcpyAsync(ptr0, sizeof(ArgsInfo), (void *)&usrArgs, sizeof(ArgsInfo), ACL_MEMCPY_HOST_TO_DEVICE, stream1);28+ error = aclrtMemcpyAsync(ptr0, sizeof(ArgsInfo), (void *)&usrArgs, sizeof(ArgsInfo), ACL_MEMCPY_HOST_TO_DEVICE, stream1);
29- 29+ 
30- // 下发任务30+ // 下发任务
31- uint32_t blockDim = 32;31+ uint32_t blockDim = 32;
32- aclrtLaunchKernelV2(funcHandle, blockDim, (void *)&usrArgs, sizeof(ArgsInfo), nullptr, stream1);32+ aclrtLaunchKernelV2(funcHandle, blockDim, (void *)&usrArgs, sizeof(ArgsInfo), nullptr, stream1);
33- 33+ 
34- // 流同步以同步释放申请内存34+ // 流同步以同步释放申请内存
35- aclrtSynchronizeStream(stream1);35+ aclrtSynchronizeStream(stream1);
36- 36+ 
37- // 释放内存，释放前需要流同步37+ // 释放内存，释放前需要流同步
38- aclrtFree(ptr0, stream1);38+ aclrtFree(ptr0, stream1);
39- ......39+ ......
40- 40+ 
41- // 流同步41+ // 流同步
42- aclrtSynchronizeStream(stream1);42+ aclrtSynchronizeStream(stream1);
43- 43+ 
44- return 0;44+ return 0;
45-}45+}
46-```46+```
47- 47+ 
48-相比之下，本章所描述的Stream有序内存分配机制将内存分配与释放操作融入到Stream调度序列中，以管理内存。这种方式将内存管理与Stream中的任务执行紧密结合，无需显式同步Stream中的任务即可进行内存管理，并且可依靠Stream本身的保序机制确保操作的有序执行。此外，Runtime还提供内存复用的能力，能够全面支持复杂的内存管理场景。48+相比之下，本章所描述的Stream有序内存分配机制将内存分配与释放操作融入到Stream调度序列中，以管理内存。这种方式将内存管理与Stream中的任务执行紧密结合，无需显式同步Stream中的任务即可进行内存管理，并且可依靠Stream本身的保序机制确保操作的有序执行。此外，Runtime还提供内存复用的能力，能够全面支持复杂的内存管理场景。
49- 49+ 
50-```cpp50+```cpp
51-#include "acl/acl_rt.h"51+#include "acl/acl_rt.h"
52-#include "acl/acl.h"52+#include "acl/acl.h"
53- 53+ 
54-int main() {54+int main() {
55- typedef struct {55+ typedef struct {
56- ......56+ ......
57- } ArgsInfo;57+ } ArgsInfo;
58- 58+ 
59- void *ptr0 = nullptr;59+ void *ptr0 = nullptr;
60- aclrtStream stream1;60+ aclrtStream stream1;
61- 61+ 
62- // 异步申请内存, testReusePool为用户创建的内存池62+ // 异步申请内存, testReusePool为用户创建的内存池
63- aclrtMemPoolMallocAsync(&ptr0, sizeof(ArgsInfo), testReusePool, stream1);63+ aclrtMemPoolMallocAsync(&ptr0, sizeof(ArgsInfo), testReusePool, stream1);
64- ......64+ ......
65- 65+ 
66- // 配置任务下发66+ // 配置任务下发
67- ArgsInfo usrArgs;67+ ArgsInfo usrArgs;
68- // 拷贝信息到device侧申请内存68+ // 拷贝信息到device侧申请内存
69- error = aclrtMemcpyAsync(ptr0, sizeof(ArgsInfo), (void *)&usrArgs, sizeof(ArgsInfo), ACL_MEMCPY_HOST_TO_DEVICE, stream1);69+ error = aclrtMemcpyAsync(ptr0, sizeof(ArgsInfo), (void *)&usrArgs, sizeof(ArgsInfo), ACL_MEMCPY_HOST_TO_DEVICE, stream1);
70- 70+ 
71- // 下发任务71+ // 下发任务
72- uint32_t blockDim = 32;72+ uint32_t blockDim = 32;
73- aclrtLaunchKernelV2(funcHandle, blockDim, (void *)&usrArgs, sizeof(ArgsInfo), nullptr, stream1);73+ aclrtLaunchKernelV2(funcHandle, blockDim, (void *)&usrArgs, sizeof(ArgsInfo), nullptr, stream1);
74- 74+ 
75- // 异步释放内存，无需进行流同步75+ // 异步释放内存，无需进行流同步
76- aclrtMemPoolFreeAsync(ptr0, stream1);76+ aclrtMemPoolFreeAsync(ptr0, stream1);
77- ......77+ ......
78- 78+ 
79- // 流同步79+ // 流同步
80- aclrtSynchronizeStream(stream1);80+ aclrtSynchronizeStream(stream1);
81- 81+ 
82- return 0;82+ return 0;
83-}83+}
84- 84+ 
85-```85+```
86- 86+ 
87-## 内存复用机制87+## 内存复用机制
88-调用aclrtMemPoolFreeAsync接口时，仅将内存归还至内存池，而不实际释放物理内存，以便后续任务能够复用这些物理内存，从而避免频繁申请和释放物理内存，提升性能。复用内存时，会根据本次任务所需的内存大小选择符合大小最接近的空闲内存。88+调用aclrtMemPoolFreeAsync接口时，仅将内存归还至内存池，而不实际释放物理内存，以便后续任务能够复用这些物理内存，从而避免频繁申请和释放物理内存，提升性能。复用内存时，会根据本次任务所需的内存大小选择符合大小最接近的空闲内存。
89- 89+ 
90-![](figures/SOMA_Reuse.png)90+![](figures/SOMA_Reuse.png)
91- 91+ 
92-当内存池中空闲的物理内存超过指定阈值（该阈值可配置，默认值为0）时，在下一次Stream同步（例如调用aclrtSynchronizeStream接口）时，系统将尝试真正释放空闲的物理内存。92+当内存池中空闲的物理内存超过指定阈值（该阈值可配置，默认值为0）时，在下一次Stream同步（例如调用aclrtSynchronizeStream接口）时，系统将尝试真正释放空闲的物理内存。
93- 93+ 
94-目前支持在一个Stream中复用内存，也支持在两个Stream之间复用内存：94+目前支持在一个Stream中复用内存，也支持在两个Stream之间复用内存：
95-* 一个Stream内进行内存复用时，基于下面的机制进行：在执行某个Stream的任务时，系统会查找该 Stream 中前序任务已归还到内存池中的内存，并复用这些内存资源，以提高内存利用率和减少内存分配的开销。95+* 一个Stream内进行内存复用时，基于下面的机制进行：在执行某个Stream的任务时，系统会查找该 Stream 中前序任务已归还到内存池中的内存，并复用这些内存资源，以提高内存利用率和减少内存分配的开销。
96-* 两个Stream之间复用内存，支持以下几种类型：96+* 两个Stream之间复用内存，支持以下几种类型：
97- * 事件依赖内存复用：在执行某个Stream的任务时，系统会查找与该Stream通过Event关联的其他Stream，并复用这些关联Stream中的任务已归还到内存池中的内存。此机制适用于用户应用程序中通过Event实现Stream间任务同步的场景。97+ * 事件依赖内存复用：在执行某个Stream的任务时，系统会查找与该Stream通过Event关联的其他Stream，并复用这些关联Stream中的任务已归还到内存池中的内存。此机制适用于用户应用程序中通过Event实现Stream间任务同步的场景。
98- * 机会主义内存复用：在执行某个Stream的任务时，系统会检索内存池中可复用的内存，但不保证内存复用一定成功。当内存复用失败时，程序会报错停止。98+ * 机会主义内存复用：在执行某个Stream的任务时，系统会检索内存池中可复用的内存，但不保证内存复用一定成功。当内存复用失败时，程序会报错停止。
99- * 隐式依赖内存复用：在执行某个Stream的任务时，系统会检索内存池中可复用的内存。若这些内存曾被其他Stream使用，但相关Stream之间不存在任务依赖关系，则系统将自动实现相关Stream之间的同步等待，以确保前一个Stream中的任务对内存的访问已经结束，从而实现安全的内存复用。99+ * 隐式依赖内存复用：在执行某个Stream的任务时，系统会检索内存池中可复用的内存。若这些内存曾被其他Stream使用，但相关Stream之间不存在任务依赖关系，则系统将自动实现相关Stream之间的同步等待，以确保前一个Stream中的任务对内存的访问已经结束，从而实现安全的内存复用。
100- 100+ 
101-## 应用场景101+## 应用场景
102-以下代码示例展现了应用异步内存申请与释放的场景，结合aclrtLaunchKernelV2接口下发任务。代码仅做参考，不能直接复制编译，需要根据实际环境和需求进行调整。102+以下代码示例展现了应用异步内存申请与释放的场景，结合aclrtLaunchKernelV2接口下发任务。代码仅做参考，不能直接复制编译，需要根据实际环境和需求进行调整。
103-```cpp103+```cpp
104-#include <stdio.h>104+#include <stdio.h>
105-#include <string.h>105+#include <string.h>
106-#include <unistd.h>106+#include <unistd.h>
107-#include <time.h>107+#include <time.h>
108-#include <gtest/gtest.h>108+#include <gtest/gtest.h>
109-#include "acl/acl_rt.h"109+#include "acl/acl_rt.h"
110-#include "acl/acl.h"110+#include "acl/acl.h"
111- 111+ 
112-int main() {112+int main() {
113- uint32_t devid = 0;113+ uint32_t devid = 0;
114- aclInit(NULL);114+ aclInit(NULL);
115- aclrtSetDevice(devid);115+ aclrtSetDevice(devid);
116- 116+ 
117- // 创建Context和Stream117+ // 创建Context和Stream
118- aclrtContext context;118+ aclrtContext context;
119- aclrtStream stream1;119+ aclrtStream stream1;
120- 120+ 
121- aclrtCreateContext(&context, 0);121+ aclrtCreateContext(&context, 0);
122- aclrtCreateStream(&stream1);122+ aclrtCreateStream(&stream1);
123- 123+ 
124- // 设置内存池属性124+ // 设置内存池属性
125- aclrtMemLocation testLoc = {125+ aclrtMemLocation testLoc = {
126- 0, // id126+ 0, // id
127- ACL_MEM_LOCATION_TYPE_DEVICE // type127+ ACL_MEM_LOCATION_TYPE_DEVICE // type
128- };128+ };
129- aclrtMemPoolProps testProp = {129+ aclrtMemPoolProps testProp = {
130- ACL_MEM_ALLOCATION_TYPE_PINNED, // allocType130+ ACL_MEM_ALLOCATION_TYPE_PINNED, // allocType
131- ACL_MEM_HANDLE_TYPE_NONE, // handleType131+ ACL_MEM_HANDLE_TYPE_NONE, // handleType
132- testLoc, // location132+ testLoc, // location
133- 14UL << 30, // maxSize = 14GB， 内存池大小为14G133+ 14UL << 30, // maxSize = 14GB， 内存池大小为14G
134- {0} // reserved134+ {0} // reserved
135- };135+ };
136- 136+ 
137- // 创建内存池137+ // 创建内存池
138- aclrtMemPool testReusePool;138+ aclrtMemPool testReusePool;
139- auto ret = aclrtMemPoolCreate(&testReusePool, &testProp);139+ auto ret = aclrtMemPoolCreate(&testReusePool, &testProp);
140- if (ret != ACL_SUCCESS) {140+ if (ret != ACL_SUCCESS) {
141- fprintf(stderr, "Failed to create memory pool\n");141+ fprintf(stderr, "Failed to create memory pool\n");
142- return -1;142+ return -1;
143- }143+ }
144- 144+ 
145- const size_t GB_TO_B = 1024ULL * 1024 * 1024;145+ const size_t GB_TO_B = 1024ULL * 1024 * 1024;
146- 146+ 
147- // 定义算子信息结构体147+ // 定义算子信息结构体
148- typedef struct {148+ typedef struct {
149- void *input_x;149+ void *input_x;
150- void *input_y;150+ void *input_y;
151- void *output_z;151+ void *output_z;
152- } ArgsInfo;152+ } ArgsInfo;
153- 153+ 
154- aclrtBinHandle bin_handle = nullptr;154+ aclrtBinHandle bin_handle = nullptr;
155- aclrtFuncHandle func_handle;155+ aclrtFuncHandle func_handle;
156- aclError aclrtBinaryGetFunction(binHandle, "add_custom", &funcHandle);156+ aclError aclrtBinaryGetFunction(binHandle, "add_custom", &funcHandle);
157- 157+ 
158- void *ptr0 = nullptr;158+ void *ptr0 = nullptr;
159- void *ptr1 = nullptr;159+ void *ptr1 = nullptr;
160- void *ptr2 = nullptr;160+ void *ptr2 = nullptr;
161- void *ptr3 = nullptr;161+ void *ptr3 = nullptr;
162- 162+ 
163- // 异步申请内存163+ // 异步申请内存
164- aclrtMemPoolMallocAsync(&ptr1, sizeof(uint64_t), testReusePool, stream1);164+ aclrtMemPoolMallocAsync(&ptr1, sizeof(uint64_t), testReusePool, stream1);
165- aclrtMemPoolMallocAsync(&ptr2, sizeof(uint64_t), testReusePool, stream1);165+ aclrtMemPoolMallocAsync(&ptr2, sizeof(uint64_t), testReusePool, stream1);
166- aclrtMemPoolMallocAsync(&ptr3, sizeof(uint64_t), testReusePool, stream1);166+ aclrtMemPoolMallocAsync(&ptr3, sizeof(uint64_t), testReusePool, stream1);
167- aclrtMemPoolMallocAsync(&ptr0, sizeof(ArgsInfo), testReusePool, stream1);167+ aclrtMemPoolMallocAsync(&ptr0, sizeof(ArgsInfo), testReusePool, stream1);
168- 168+ 
169- // 配置任务下发169+ // 配置任务下发
170- ArgsInfo usrArgs;170+ ArgsInfo usrArgs;
171- usrAgrs.input_x = ptr1171+ usrAgrs.input_x = ptr1
172- usrAgrs.input_y = ptr2;172+ usrAgrs.input_y = ptr2;
173- usrAgrs.output_z = ptr3;173+ usrAgrs.output_z = ptr3;
174- 174+ 
175- error = aclrtMemcpyAsync(devPtr, sizeof(ArgsInfo), (void *)&usrArgs, sizeof(ArgsInfo), ACL_MEMCPY_HOST_TO_DEVICE, stream1);175+ error = aclrtMemcpyAsync(devPtr, sizeof(ArgsInfo), (void *)&usrArgs, sizeof(ArgsInfo), ACL_MEMCPY_HOST_TO_DEVICE, stream1);
176- 176+ 
177- // 下发任务177+ // 下发任务
178- uint32_t blockDim = 32;178+ uint32_t blockDim = 32;
179- aclrtLaunchKernelV2(funcHandle, blockDim, (void *)&usrArgs, sizeof(ArgsInfo), nullptr, stream1);179+ aclrtLaunchKernelV2(funcHandle, blockDim, (void *)&usrArgs, sizeof(ArgsInfo), nullptr, stream1);
180- 180+ 
181- // 异步释放内存，此前无需进行流同步181+ // 异步释放内存，此前无需进行流同步
182- aclrtMemPoolFreeAsync(ptr0, stream1);182+ aclrtMemPoolFreeAsync(ptr0, stream1);
183- aclrtMemPoolFreeAsync(ptr1, stream1);183+ aclrtMemPoolFreeAsync(ptr1, stream1);
184- aclrtMemPoolFreeAsync(ptr2, stream1);184+ aclrtMemPoolFreeAsync(ptr2, stream1);
185- aclrtMemPoolFreeAsync(ptr3, stream1);185+ aclrtMemPoolFreeAsync(ptr3, stream1);
186- 186+ 
187- // 流同步187+ // 流同步
188- aclrtSynchronizeStream(stream1);188+ aclrtSynchronizeStream(stream1);
189- 189+ 
190- // 销毁内存池、Stream和Context190+ // 销毁内存池、Stream和Context
191- aclrtMemPoolDestroy(testReusePool);191+ aclrtMemPoolDestroy(testReusePool);
192- aclrtDestroyStream(stream1);192+ aclrtDestroyStream(stream1);
193- aclrtDestroyContext(context);193+ aclrtDestroyContext(context);
194- 194+ 
195- aclrtResetDevice(devid);195+ aclrtResetDevice(devid);
196- aclFinalize();196+ aclFinalize();
197- return 0;197+ return 0;
198-}198+}
199-```199+```
Rdocs/api_docs/Stream有序内存分配.md→docs/03_api_ref/Stream有序内存分配.md+0-0
文件重命名但无更改。
Rdocs/api_docs/Stream管理.md→docs/03_api_ref/Stream管理.md+0-0
文件重命名但无更改。
Rdocs/api_docs/Tensor数据传输.md→docs/03_api_ref/Tensor数据传输.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclAppLog.md→docs/03_api_ref/aclAppLog.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclCANNPackageName.md→docs/03_api_ref/aclCANNPackageName.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclCANNPackageVersion.md→docs/03_api_ref/aclCANNPackageVersion.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclCannAttr.md→docs/03_api_ref/aclCannAttr.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclCreateDataBuffer.md→docs/03_api_ref/aclCreateDataBuffer.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclDataBuffer.md→docs/03_api_ref/aclDataBuffer.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclDataType.md→docs/03_api_ref/aclDataType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclDataTypeSize.md→docs/03_api_ref/aclDataTypeSize.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclDestroyDataBuffer.md→docs/03_api_ref/aclDestroyDataBuffer.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclDeviceInfo.md→docs/03_api_ref/aclDeviceInfo.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclError.md→docs/03_api_ref/aclError.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclFinalize.md→docs/03_api_ref/aclFinalize.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclFinalizeCallbackRegister.md→docs/03_api_ref/aclFinalizeCallbackRegister.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclFinalizeCallbackUnRegister.md→docs/03_api_ref/aclFinalizeCallbackUnRegister.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclFinalizeReference.md→docs/03_api_ref/aclFinalizeReference.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclFloat16.md→docs/03_api_ref/aclFloat16.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclFloat16ToFloat.md→docs/03_api_ref/aclFloat16ToFloat.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclFloatToFloat16.md→docs/03_api_ref/aclFloatToFloat16.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclFormat.md→docs/03_api_ref/aclFormat.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclGetCannAttribute.md→docs/03_api_ref/aclGetCannAttribute.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclGetCannAttributeList.md→docs/03_api_ref/aclGetCannAttributeList.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclGetDataBufferAddr.md→docs/03_api_ref/aclGetDataBufferAddr.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclGetDataBufferSizeV2.md→docs/03_api_ref/aclGetDataBufferSizeV2.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclGetDataBufferSize（废弃）.md→docs/03_api_ref/aclGetDataBufferSize（废弃）.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclGetDeviceCapability.md→docs/03_api_ref/aclGetDeviceCapability.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclGetRecentErrMsg.md→docs/03_api_ref/aclGetRecentErrMsg.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclInit.md→docs/03_api_ref/aclInit.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclInitCallbackRegister.md→docs/03_api_ref/aclInitCallbackRegister.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclInitCallbackUnRegister.md→docs/03_api_ref/aclInitCallbackUnRegister.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclMemType.md→docs/03_api_ref/aclMemType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclRecoverAllHcclTasks.md→docs/03_api_ref/aclRecoverAllHcclTasks.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclRegisterCallbackType.md→docs/03_api_ref/aclRegisterCallbackType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclSysParamOpt.md→docs/03_api_ref/aclSysParamOpt.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclUpdateDataBuffer.md→docs/03_api_ref/aclUpdateDataBuffer.md+0-0
文件重命名但无更改。
Rdocs/api_docs/acldumpGetPath.md→docs/03_api_ref/acldumpGetPath.md+0-0
文件重命名但无更改。
Rdocs/api_docs/acldumpRegCallback.md→docs/03_api_ref/acldumpRegCallback.md+0-0
文件重命名但无更改。
Rdocs/api_docs/acldumpType.md→docs/03_api_ref/acldumpType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/acldumpUnregCallback.md→docs/03_api_ref/acldumpUnregCallback.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlFinalizeDump.md→docs/03_api_ref/aclmdlFinalizeDump.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlInitDump.md→docs/03_api_ref/aclmdlInitDump.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRI.md→docs/03_api_ref/aclmdlRI.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIAbort.md→docs/03_api_ref/aclmdlRIAbort.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIBindStream.md→docs/03_api_ref/aclmdlRIBindStream.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIBuildBegin.md→docs/03_api_ref/aclmdlRIBuildBegin.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIBuildEnd.md→docs/03_api_ref/aclmdlRIBuildEnd.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRICaptureBegin.md→docs/03_api_ref/aclmdlRICaptureBegin.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRICaptureEnd.md→docs/03_api_ref/aclmdlRICaptureEnd.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRICaptureGetInfo.md→docs/03_api_ref/aclmdlRICaptureGetInfo.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRICaptureMode.md→docs/03_api_ref/aclmdlRICaptureMode.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRICaptureStatus.md→docs/03_api_ref/aclmdlRICaptureStatus.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRICaptureTaskGrpBegin.md→docs/03_api_ref/aclmdlRICaptureTaskGrpBegin.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRICaptureTaskGrpEnd.md→docs/03_api_ref/aclmdlRICaptureTaskGrpEnd.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRICaptureTaskUpdateBegin.md→docs/03_api_ref/aclmdlRICaptureTaskUpdateBegin.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRICaptureTaskUpdateEnd.md→docs/03_api_ref/aclmdlRICaptureTaskUpdateEnd.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRICaptureThreadExchangeMode.md→docs/03_api_ref/aclmdlRICaptureThreadExchangeMode.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIDebugJsonPrint.md→docs/03_api_ref/aclmdlRIDebugJsonPrint.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIDebugPrint.md→docs/03_api_ref/aclmdlRIDebugPrint.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIDestroy.md→docs/03_api_ref/aclmdlRIDestroy.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIDestroyRegisterCallback.md→docs/03_api_ref/aclmdlRIDestroyRegisterCallback.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIDestroyUnregisterCallback.md→docs/03_api_ref/aclmdlRIDestroyUnregisterCallback.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIEndTask.md→docs/03_api_ref/aclmdlRIEndTask.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIExecute.md→docs/03_api_ref/aclmdlRIExecute.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIExecuteAsync.md→docs/03_api_ref/aclmdlRIExecuteAsync.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIGetName.md→docs/03_api_ref/aclmdlRIGetName.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIGetStreams.md→docs/03_api_ref/aclmdlRIGetStreams.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIGetTasksByStream.md→docs/03_api_ref/aclmdlRIGetTasksByStream.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRISetName.md→docs/03_api_ref/aclmdlRISetName.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRITask.md→docs/03_api_ref/aclmdlRITask.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRITaskGetSeqId.md→docs/03_api_ref/aclmdlRITaskGetSeqId.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRITaskGetType.md→docs/03_api_ref/aclmdlRITaskGetType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRITaskType.md→docs/03_api_ref/aclmdlRITaskType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlRIUnbindStream.md→docs/03_api_ref/aclmdlRIUnbindStream.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclmdlSetDump.md→docs/03_api_ref/aclmdlSetDump.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclopStartDumpArgs.md→docs/03_api_ref/aclopStartDumpArgs.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclopStopDumpArgs.md→docs/03_api_ref/aclopStopDumpArgs.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofAicoreMetrics.md→docs/03_api_ref/aclprofAicoreMetrics.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofConfig.md→docs/03_api_ref/aclprofConfig.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofCreateConfig.md→docs/03_api_ref/aclprofCreateConfig.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofCreateStamp.md→docs/03_api_ref/aclprofCreateStamp.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofCreateStepInfo.md→docs/03_api_ref/aclprofCreateStepInfo.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofCreateSubscribeConfig.md→docs/03_api_ref/aclprofCreateSubscribeConfig.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofDestroyConfig.md→docs/03_api_ref/aclprofDestroyConfig.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofDestroyStamp.md→docs/03_api_ref/aclprofDestroyStamp.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofDestroyStepInfo.md→docs/03_api_ref/aclprofDestroyStepInfo.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofDestroySubscribeConfig.md→docs/03_api_ref/aclprofDestroySubscribeConfig.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofFinalize.md→docs/03_api_ref/aclprofFinalize.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofGetModelId.md→docs/03_api_ref/aclprofGetModelId.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofGetOpDescSize.md→docs/03_api_ref/aclprofGetOpDescSize.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofGetOpDuration.md→docs/03_api_ref/aclprofGetOpDuration.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofGetOpEnd.md→docs/03_api_ref/aclprofGetOpEnd.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofGetOpName.md→docs/03_api_ref/aclprofGetOpName.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofGetOpNameLen.md→docs/03_api_ref/aclprofGetOpNameLen.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofGetOpNum.md→docs/03_api_ref/aclprofGetOpNum.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofGetOpStart.md→docs/03_api_ref/aclprofGetOpStart.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofGetOpType.md→docs/03_api_ref/aclprofGetOpType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofGetOpTypeLen.md→docs/03_api_ref/aclprofGetOpTypeLen.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofGetStepTimestamp.md→docs/03_api_ref/aclprofGetStepTimestamp.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofInit.md→docs/03_api_ref/aclprofInit.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofMark.md→docs/03_api_ref/aclprofMark.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofMarkEx.md→docs/03_api_ref/aclprofMarkEx.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofModelSubscribe.md→docs/03_api_ref/aclprofModelSubscribe.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofModelUnSubscribe.md→docs/03_api_ref/aclprofModelUnSubscribe.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofPop.md→docs/03_api_ref/aclprofPop.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofPush.md→docs/03_api_ref/aclprofPush.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofRangeStart.md→docs/03_api_ref/aclprofRangeStart.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofRangeStop.md→docs/03_api_ref/aclprofRangeStop.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofSetConfig.md→docs/03_api_ref/aclprofSetConfig.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofSetStampTraceMessage.md→docs/03_api_ref/aclprofSetStampTraceMessage.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofStart.md→docs/03_api_ref/aclprofStart.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofStepInfo.md→docs/03_api_ref/aclprofStepInfo.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofStepTag.md→docs/03_api_ref/aclprofStepTag.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofStop.md→docs/03_api_ref/aclprofStop.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclprofSubscribeConfig.md→docs/03_api_ref/aclprofSubscribeConfig.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtActiveStream.md→docs/03_api_ref/aclrtActiveStream.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAicAivTaskUpdateAttr.md→docs/03_api_ref/aclrtAicAivTaskUpdateAttr.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocBuf.md→docs/03_api_ref/aclrtAllocBuf.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocator.md→docs/03_api_ref/aclrtAllocator.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocatorAddr.md→docs/03_api_ref/aclrtAllocatorAddr.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocatorBlock.md→docs/03_api_ref/aclrtAllocatorBlock.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocatorCreateDesc.md→docs/03_api_ref/aclrtAllocatorCreateDesc.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocatorDesc.md→docs/03_api_ref/aclrtAllocatorDesc.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocatorDestroyDesc.md→docs/03_api_ref/aclrtAllocatorDestroyDesc.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocatorGetByStream.md→docs/03_api_ref/aclrtAllocatorGetByStream.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocatorRegister.md→docs/03_api_ref/aclrtAllocatorRegister.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocatorSetAllocAdviseFuncToDesc.md→docs/03_api_ref/aclrtAllocatorSetAllocAdviseFuncToDesc.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocatorSetAllocFuncToDesc.md→docs/03_api_ref/aclrtAllocatorSetAllocFuncToDesc.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocatorSetFreeFuncToDesc.md→docs/03_api_ref/aclrtAllocatorSetFreeFuncToDesc.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocatorSetGetAddrFromBlockFuncToDesc.md→docs/03_api_ref/aclrtAllocatorSetGetAddrFromBlockFuncToDesc.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocatorSetObjToDesc.md→docs/03_api_ref/aclrtAllocatorSetObjToDesc.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAllocatorUnregister.md→docs/03_api_ref/aclrtAllocatorUnregister.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtAppendBufChain.md→docs/03_api_ref/aclrtAppendBufChain.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtArgsHandle.md→docs/03_api_ref/aclrtArgsHandle.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBarrierCmoInfo.md→docs/03_api_ref/aclrtBarrierCmoInfo.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBarrierTaskInfo.md→docs/03_api_ref/aclrtBarrierTaskInfo.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBinHandle.md→docs/03_api_ref/aclrtBinHandle.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBinaryGetDevAddress.md→docs/03_api_ref/aclrtBinaryGetDevAddress.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBinaryGetFunction.md→docs/03_api_ref/aclrtBinaryGetFunction.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBinaryGetFunctionByEntry.md→docs/03_api_ref/aclrtBinaryGetFunctionByEntry.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBinaryLoad.md→docs/03_api_ref/aclrtBinaryLoad.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBinaryLoadFromData.md→docs/03_api_ref/aclrtBinaryLoadFromData.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBinaryLoadFromFile.md→docs/03_api_ref/aclrtBinaryLoadFromFile.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBinaryLoadOption.md→docs/03_api_ref/aclrtBinaryLoadOption.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBinaryLoadOptionType.md→docs/03_api_ref/aclrtBinaryLoadOptionType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBinaryLoadOptionValue.md→docs/03_api_ref/aclrtBinaryLoadOptionValue.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBinaryLoadOptions.md→docs/03_api_ref/aclrtBinaryLoadOptions.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBinarySetExceptionCallback.md→docs/03_api_ref/aclrtBinarySetExceptionCallback.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtBinaryUnLoad.md→docs/03_api_ref/aclrtBinaryUnLoad.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCacheLastTaskOpInfo.md→docs/03_api_ref/aclrtCacheLastTaskOpInfo.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCheckArchCompatibility.md→docs/03_api_ref/aclrtCheckArchCompatibility.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCheckMemType.md→docs/03_api_ref/aclrtCheckMemType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCmoAsync.md→docs/03_api_ref/aclrtCmoAsync.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCmoAsyncWithBarrier.md→docs/03_api_ref/aclrtCmoAsyncWithBarrier.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCmoAsyncWithDesc.md→docs/03_api_ref/aclrtCmoAsyncWithDesc.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCmoGetDescSize.md→docs/03_api_ref/aclrtCmoGetDescSize.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCmoSetDesc.md→docs/03_api_ref/aclrtCmoSetDesc.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCmoType.md→docs/03_api_ref/aclrtCmoType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCmoWaitBarrier.md→docs/03_api_ref/aclrtCmoWaitBarrier.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCntNotify.md→docs/03_api_ref/aclrtCntNotify.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCntNotifyCreate.md→docs/03_api_ref/aclrtCntNotifyCreate.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCntNotifyDestroy.md→docs/03_api_ref/aclrtCntNotifyDestroy.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCntNotifyGetId.md→docs/03_api_ref/aclrtCntNotifyGetId.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCntNotifyRecord.md→docs/03_api_ref/aclrtCntNotifyRecord.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCntNotifyRecordInfo.md→docs/03_api_ref/aclrtCntNotifyRecordInfo.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCntNotifyRecordMode.md→docs/03_api_ref/aclrtCntNotifyRecordMode.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCntNotifyReset.md→docs/03_api_ref/aclrtCntNotifyReset.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCntNotifyWaitInfo.md→docs/03_api_ref/aclrtCntNotifyWaitInfo.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCntNotifyWaitMode.md→docs/03_api_ref/aclrtCntNotifyWaitMode.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCntNotifyWaitWithTimeout.md→docs/03_api_ref/aclrtCntNotifyWaitWithTimeout.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCompareDataType.md→docs/03_api_ref/aclrtCompareDataType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCondition.md→docs/03_api_ref/aclrtCondition.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtContext.md→docs/03_api_ref/aclrtContext.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCopyBufRef.md→docs/03_api_ref/aclrtCopyBufRef.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCreateBinary.md→docs/03_api_ref/aclrtCreateBinary.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCreateContext.md→docs/03_api_ref/aclrtCreateContext.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCreateEvent.md→docs/03_api_ref/aclrtCreateEvent.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCreateEventExWithFlag.md→docs/03_api_ref/aclrtCreateEventExWithFlag.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCreateEventWithFlag.md→docs/03_api_ref/aclrtCreateEventWithFlag.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCreateGroupInfo.md→docs/03_api_ref/aclrtCreateGroupInfo.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCreateLabel.md→docs/03_api_ref/aclrtCreateLabel.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCreateLabelList.md→docs/03_api_ref/aclrtCreateLabelList.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCreateNotify.md→docs/03_api_ref/aclrtCreateNotify.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCreateStream.md→docs/03_api_ref/aclrtCreateStream.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCreateStreamConfigHandle.md→docs/03_api_ref/aclrtCreateStreamConfigHandle.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCreateStreamV2.md→docs/03_api_ref/aclrtCreateStreamV2.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCreateStreamWithConfig.md→docs/03_api_ref/aclrtCreateStreamWithConfig.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCtxGetCurrentDefaultStream.md→docs/03_api_ref/aclrtCtxGetCurrentDefaultStream.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCtxGetFloatOverflowAddr.md→docs/03_api_ref/aclrtCtxGetFloatOverflowAddr.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCtxGetSysParamOpt.md→docs/03_api_ref/aclrtCtxGetSysParamOpt.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtCtxSetSysParamOpt.md→docs/03_api_ref/aclrtCtxSetSysParamOpt.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDestroyBinary.md→docs/03_api_ref/aclrtDestroyBinary.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDestroyContext.md→docs/03_api_ref/aclrtDestroyContext.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDestroyEvent.md→docs/03_api_ref/aclrtDestroyEvent.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDestroyGroupInfo.md→docs/03_api_ref/aclrtDestroyGroupInfo.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDestroyLabel.md→docs/03_api_ref/aclrtDestroyLabel.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDestroyLabelList.md→docs/03_api_ref/aclrtDestroyLabelList.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDestroyNotify.md→docs/03_api_ref/aclrtDestroyNotify.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDestroyStream.md→docs/03_api_ref/aclrtDestroyStream.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDestroyStreamConfigHandle.md→docs/03_api_ref/aclrtDestroyStreamConfigHandle.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDestroyStreamForce.md→docs/03_api_ref/aclrtDestroyStreamForce.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDevAttr.md→docs/03_api_ref/aclrtDevAttr.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDevFeatureType.md→docs/03_api_ref/aclrtDevFeatureType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDevResLimitType.md→docs/03_api_ref/aclrtDevResLimitType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDeviceCanAccessPeer.md→docs/03_api_ref/aclrtDeviceCanAccessPeer.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDeviceDisablePeerAccess.md→docs/03_api_ref/aclrtDeviceDisablePeerAccess.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDeviceEnablePeerAccess.md→docs/03_api_ref/aclrtDeviceEnablePeerAccess.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDeviceGetBareTgid.md→docs/03_api_ref/aclrtDeviceGetBareTgid.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDeviceGetStreamPriorityRange.md→docs/03_api_ref/aclrtDeviceGetStreamPriorityRange.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDeviceGetUuid.md→docs/03_api_ref/aclrtDeviceGetUuid.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDevicePeerAccessStatus.md→docs/03_api_ref/aclrtDevicePeerAccessStatus.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDeviceStatus.md→docs/03_api_ref/aclrtDeviceStatus.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDeviceTaskAbort.md→docs/03_api_ref/aclrtDeviceTaskAbort.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDropoutBitMaskInfo.md→docs/03_api_ref/aclrtDropoutBitMaskInfo.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtDrvMemHandle.md→docs/03_api_ref/aclrtDrvMemHandle.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtEngineType.md→docs/03_api_ref/aclrtEngineType.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtEvent.md→docs/03_api_ref/aclrtEvent.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtEventElapsedTime.md→docs/03_api_ref/aclrtEventElapsedTime.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtEventGetTimestamp.md→docs/03_api_ref/aclrtEventGetTimestamp.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtEventRecordedStatus.md→docs/03_api_ref/aclrtEventRecordedStatus.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtEventStatus.md→docs/03_api_ref/aclrtEventStatus.md+0-0
文件重命名但无更改。
Rdocs/api_docs/aclrtEventWaitStatus.md→docs/03_api_ref/aclrtEventWaitStatus.md+0-0
Rdocs/api_docs/aclrtFloatOverflowMode.md→docs/03_api_ref/aclrtFloatOverflowMode.md+0-0
Rdocs/api_docs/aclrtFree.md→docs/03_api_ref/aclrtFree.md+0-0
Rdocs/api_docs/aclrtFreeBuf.md→docs/03_api_ref/aclrtFreeBuf.md+0-0
Rdocs/api_docs/aclrtFreeHost.md→docs/03_api_ref/aclrtFreeHost.md+0-0
Rdocs/api_docs/aclrtFreeHostWithDevSync.md→docs/03_api_ref/aclrtFreeHostWithDevSync.md+0-0
Rdocs/api_docs/aclrtFreePhysical.md→docs/03_api_ref/aclrtFreePhysical.md+0-0
Rdocs/api_docs/aclrtFreeWithDevSync.md→docs/03_api_ref/aclrtFreeWithDevSync.md+0-0
Rdocs/api_docs/aclrtFuncAttribute.md→docs/03_api_ref/aclrtFuncAttribute.md+0-0
Rdocs/api_docs/aclrtFuncHandle.md→docs/03_api_ref/aclrtFuncHandle.md+0-0
Rdocs/api_docs/aclrtGetAllGroupInfo.md→docs/03_api_ref/aclrtGetAllGroupInfo.md+0-0
Rdocs/api_docs/aclrtGetArgsFromExceptionInfo.md→docs/03_api_ref/aclrtGetArgsFromExceptionInfo.md+0-0
Rdocs/api_docs/aclrtGetBufChainNum.md→docs/03_api_ref/aclrtGetBufChainNum.md+0-0
Rdocs/api_docs/aclrtGetBufData.md→docs/03_api_ref/aclrtGetBufData.md+0-0
Rdocs/api_docs/aclrtGetBufDataLen.md→docs/03_api_ref/aclrtGetBufDataLen.md+0-0
Rdocs/api_docs/aclrtGetBufFromChain.md→docs/03_api_ref/aclrtGetBufFromChain.md+0-0
Rdocs/api_docs/aclrtGetBufUserData.md→docs/03_api_ref/aclrtGetBufUserData.md+0-0
Rdocs/api_docs/aclrtGetCurrentContext.md→docs/03_api_ref/aclrtGetCurrentContext.md+0-0
Rdocs/api_docs/aclrtGetDevice.md→docs/03_api_ref/aclrtGetDevice.md+0-0
Rdocs/api_docs/aclrtGetDeviceCapability.md→docs/03_api_ref/aclrtGetDeviceCapability.md+0-0
Rdocs/api_docs/aclrtGetDeviceCount.md→docs/03_api_ref/aclrtGetDeviceCount.md+0-0
Rdocs/api_docs/aclrtGetDeviceIdFromExceptionInfo.md→docs/03_api_ref/aclrtGetDeviceIdFromExceptionInfo.md+0-0
Rdocs/api_docs/aclrtGetDeviceInfo.md→docs/03_api_ref/aclrtGetDeviceInfo.md+0-0
Rdocs/api_docs/aclrtGetDeviceResLimit.md→docs/03_api_ref/aclrtGetDeviceResLimit.md+0-0
Rdocs/api_docs/aclrtGetDeviceSatMode.md→docs/03_api_ref/aclrtGetDeviceSatMode.md+0-0
Rdocs/api_docs/aclrtGetDeviceUtilizationRate.md→docs/03_api_ref/aclrtGetDeviceUtilizationRate.md+0-0
Rdocs/api_docs/aclrtGetDevicesTopo.md→docs/03_api_ref/aclrtGetDevicesTopo.md+0-0
Rdocs/api_docs/aclrtGetErrorCodeFromExceptionInfo.md→docs/03_api_ref/aclrtGetErrorCodeFromExceptionInfo.md+0-0
Rdocs/api_docs/aclrtGetErrorVerbose.md→docs/03_api_ref/aclrtGetErrorVerbose.md+0-0
Rdocs/api_docs/aclrtGetEventAvailNum.md→docs/03_api_ref/aclrtGetEventAvailNum.md+0-0
Rdocs/api_docs/aclrtGetEventId.md→docs/03_api_ref/aclrtGetEventId.md+0-0
Rdocs/api_docs/aclrtGetFuncHandleFromExceptionInfo.md→docs/03_api_ref/aclrtGetFuncHandleFromExceptionInfo.md+0-0
Rdocs/api_docs/aclrtGetFunctionAddr.md→docs/03_api_ref/aclrtGetFunctionAddr.md+0-0
Rdocs/api_docs/aclrtGetFunctionAttribute.md→docs/03_api_ref/aclrtGetFunctionAttribute.md+0-0
Rdocs/api_docs/aclrtGetFunctionName.md→docs/03_api_ref/aclrtGetFunctionName.md+0-0
Rdocs/api_docs/aclrtGetFunctionSize.md→docs/03_api_ref/aclrtGetFunctionSize.md+0-0
Rdocs/api_docs/aclrtGetGroupCount.md→docs/03_api_ref/aclrtGetGroupCount.md+0-0
Rdocs/api_docs/aclrtGetGroupInfoDetail.md→docs/03_api_ref/aclrtGetGroupInfoDetail.md+0-0
Rdocs/api_docs/aclrtGetHardwareSyncAddr.md→docs/03_api_ref/aclrtGetHardwareSyncAddr.md+0-0
Rdocs/api_docs/aclrtGetLastError.md→docs/03_api_ref/aclrtGetLastError.md+0-0
Rdocs/api_docs/aclrtGetLogicDevIdByPhyDevId.md→docs/03_api_ref/aclrtGetLogicDevIdByPhyDevId.md+0-0
Rdocs/api_docs/aclrtGetLogicDevIdByUserDevId.md→docs/03_api_ref/aclrtGetLogicDevIdByUserDevId.md+0-0
Rdocs/api_docs/aclrtGetMemInfo.md→docs/03_api_ref/aclrtGetMemInfo.md+0-0
Rdocs/api_docs/aclrtGetMemUceInfo.md→docs/03_api_ref/aclrtGetMemUceInfo.md+0-0
Rdocs/api_docs/aclrtGetMemcpyDescSize.md→docs/03_api_ref/aclrtGetMemcpyDescSize.md+0-0
Rdocs/api_docs/aclrtGetNotifyId.md→docs/03_api_ref/aclrtGetNotifyId.md+0-0
Rdocs/api_docs/aclrtGetOpExecuteTimeOut.md→docs/03_api_ref/aclrtGetOpExecuteTimeOut.md+0-0
Rdocs/api_docs/aclrtGetOpTimeoutInterval.md→docs/03_api_ref/aclrtGetOpTimeoutInterval.md+0-0
Rdocs/api_docs/aclrtGetOverflowStatus.md→docs/03_api_ref/aclrtGetOverflowStatus.md+0-0
Rdocs/api_docs/aclrtGetPhyDevIdByLogicDevId.md→docs/03_api_ref/aclrtGetPhyDevIdByLogicDevId.md+0-0
Rdocs/api_docs/aclrtGetPrimaryCtxState.md→docs/03_api_ref/aclrtGetPrimaryCtxState.md+0-0
Rdocs/api_docs/aclrtGetResInCurrentThread.md→docs/03_api_ref/aclrtGetResInCurrentThread.md+0-0
Rdocs/api_docs/aclrtGetRunMode.md→docs/03_api_ref/aclrtGetRunMode.md+0-0
Rdocs/api_docs/aclrtGetSocName.md→docs/03_api_ref/aclrtGetSocName.md+0-0
Rdocs/api_docs/aclrtGetStreamAttribute.md→docs/03_api_ref/aclrtGetStreamAttribute.md+0-0
Rdocs/api_docs/aclrtGetStreamAvailableNum.md→docs/03_api_ref/aclrtGetStreamAvailableNum.md+0-0
Rdocs/api_docs/aclrtGetStreamIdFromExceptionInfo.md→docs/03_api_ref/aclrtGetStreamIdFromExceptionInfo.md+0-0
Rdocs/api_docs/aclrtGetStreamOverflowSwitch.md→docs/03_api_ref/aclrtGetStreamOverflowSwitch.md+0-0
Rdocs/api_docs/aclrtGetStreamResLimit.md→docs/03_api_ref/aclrtGetStreamResLimit.md+0-0
Rdocs/api_docs/aclrtGetSysParamOpt.md→docs/03_api_ref/aclrtGetSysParamOpt.md+0-0
Rdocs/api_docs/aclrtGetTaskIdFromExceptionInfo.md→docs/03_api_ref/aclrtGetTaskIdFromExceptionInfo.md+0-0
Rdocs/api_docs/aclrtGetThreadIdFromExceptionInfo.md→docs/03_api_ref/aclrtGetThreadIdFromExceptionInfo.md+0-0
Rdocs/api_docs/aclrtGetThreadLastTaskId.md→docs/03_api_ref/aclrtGetThreadLastTaskId.md+0-0
Rdocs/api_docs/aclrtGetUserDevIdByLogicDevId.md→docs/03_api_ref/aclrtGetUserDevIdByLogicDevId.md+0-0
Rdocs/api_docs/aclrtGetVersion.md→docs/03_api_ref/aclrtGetVersion.md+0-0
Rdocs/api_docs/aclrtGroupAttr.md→docs/03_api_ref/aclrtGroupAttr.md+0-0
Rdocs/api_docs/aclrtHacType.md→docs/03_api_ref/aclrtHacType.md+0-0
Rdocs/api_docs/aclrtHostGetDevicePointer.md→docs/03_api_ref/aclrtHostGetDevicePointer.md+0-0
Rdocs/api_docs/aclrtHostMemMapCapabilities.md→docs/03_api_ref/aclrtHostMemMapCapabilities.md+0-0
Rdocs/api_docs/aclrtHostMemMapCapability.md→docs/03_api_ref/aclrtHostMemMapCapability.md+0-0
Rdocs/api_docs/aclrtHostRegister.md→docs/03_api_ref/aclrtHostRegister.md+0-0
Rdocs/api_docs/aclrtHostRegisterType.md→docs/03_api_ref/aclrtHostRegisterType.md+0-0
Rdocs/api_docs/aclrtHostRegisterV2.md→docs/03_api_ref/aclrtHostRegisterV2.md+0-0
Rdocs/api_docs/aclrtHostUnregister.md→docs/03_api_ref/aclrtHostUnregister.md+0-0
Rdocs/api_docs/aclrtIpcEventHandle.md→docs/03_api_ref/aclrtIpcEventHandle.md+0-0
Rdocs/api_docs/aclrtIpcGetEventHandle.md→docs/03_api_ref/aclrtIpcGetEventHandle.md+0-0
Rdocs/api_docs/aclrtIpcMemAttrType.md→docs/03_api_ref/aclrtIpcMemAttrType.md+0-0
Rdocs/api_docs/aclrtIpcMemClose.md→docs/03_api_ref/aclrtIpcMemClose.md+0-0
Rdocs/api_docs/aclrtIpcMemGetExportKey.md→docs/03_api_ref/aclrtIpcMemGetExportKey.md+0-0
Rdocs/api_docs/aclrtIpcMemImportByKey.md→docs/03_api_ref/aclrtIpcMemImportByKey.md+0-0
Rdocs/api_docs/aclrtIpcMemImportPidInterServer.md→docs/03_api_ref/aclrtIpcMemImportPidInterServer.md+0-0
Rdocs/api_docs/aclrtIpcMemSetAttr.md→docs/03_api_ref/aclrtIpcMemSetAttr.md+0-0
Rdocs/api_docs/aclrtIpcMemSetImportPid.md→docs/03_api_ref/aclrtIpcMemSetImportPid.md+0-0
Rdocs/api_docs/aclrtIpcOpenEventHandle.md→docs/03_api_ref/aclrtIpcOpenEventHandle.md+0-0
Rdocs/api_docs/aclrtKernelArgsAppend.md→docs/03_api_ref/aclrtKernelArgsAppend.md+0-0
Rdocs/api_docs/aclrtKernelArgsAppendPlaceHolder.md→docs/03_api_ref/aclrtKernelArgsAppendPlaceHolder.md+0-0
Rdocs/api_docs/aclrtKernelArgsFinalize.md→docs/03_api_ref/aclrtKernelArgsFinalize.md+0-0
Rdocs/api_docs/aclrtKernelArgsGetHandleMemSize.md→docs/03_api_ref/aclrtKernelArgsGetHandleMemSize.md+0-0
Rdocs/api_docs/aclrtKernelArgsGetMemSize.md→docs/03_api_ref/aclrtKernelArgsGetMemSize.md+0-0
Rdocs/api_docs/aclrtKernelArgsGetPlaceHolderBuffer.md→docs/03_api_ref/aclrtKernelArgsGetPlaceHolderBuffer.md+0-0
Rdocs/api_docs/aclrtKernelArgsInit.md→docs/03_api_ref/aclrtKernelArgsInit.md+0-0
Rdocs/api_docs/aclrtKernelArgsInitByUserMem.md→docs/03_api_ref/aclrtKernelArgsInitByUserMem.md+0-0
Rdocs/api_docs/aclrtKernelArgsParaUpdate.md→docs/03_api_ref/aclrtKernelArgsParaUpdate.md+0-0
Rdocs/api_docs/aclrtKernelType.md→docs/03_api_ref/aclrtKernelType.md+0-0
Rdocs/api_docs/aclrtLabel.md→docs/03_api_ref/aclrtLabel.md+0-0
Rdocs/api_docs/aclrtLabelList.md→docs/03_api_ref/aclrtLabelList.md+0-0
Rdocs/api_docs/aclrtLastErrLevel.md→docs/03_api_ref/aclrtLastErrLevel.md+0-0
Rdocs/api_docs/aclrtLaunchCallback.md→docs/03_api_ref/aclrtLaunchCallback.md+0-0
Rdocs/api_docs/aclrtLaunchHostFunc.md→docs/03_api_ref/aclrtLaunchHostFunc.md+0-0
Rdocs/api_docs/aclrtLaunchKernel.md→docs/03_api_ref/aclrtLaunchKernel.md+0-0
Rdocs/api_docs/aclrtLaunchKernelAttr.md→docs/03_api_ref/aclrtLaunchKernelAttr.md+0-0
Rdocs/api_docs/aclrtLaunchKernelAttrId.md→docs/03_api_ref/aclrtLaunchKernelAttrId.md+0-0
Rdocs/api_docs/aclrtLaunchKernelAttrValue.md→docs/03_api_ref/aclrtLaunchKernelAttrValue.md+0-0
Rdocs/api_docs/aclrtLaunchKernelCfg.md→docs/03_api_ref/aclrtLaunchKernelCfg.md+0-0
Rdocs/api_docs/aclrtLaunchKernelV2.md→docs/03_api_ref/aclrtLaunchKernelV2.md+0-0
Rdocs/api_docs/aclrtLaunchKernelWithConfig.md→docs/03_api_ref/aclrtLaunchKernelWithConfig.md+0-0
Rdocs/api_docs/aclrtLaunchKernelWithHostArgs.md→docs/03_api_ref/aclrtLaunchKernelWithHostArgs.md+0-0
Rdocs/api_docs/aclrtMalloc.md→docs/03_api_ref/aclrtMalloc.md+0-0
Rdocs/api_docs/aclrtMallocAlign32.md→docs/03_api_ref/aclrtMallocAlign32.md+0-0
Rdocs/api_docs/aclrtMallocAttrType.md→docs/03_api_ref/aclrtMallocAttrType.md+0-0
Rdocs/api_docs/aclrtMallocAttrValue.md→docs/03_api_ref/aclrtMallocAttrValue.md+0-0
Rdocs/api_docs/aclrtMallocAttribute.md→docs/03_api_ref/aclrtMallocAttribute.md+0-0
Rdocs/api_docs/aclrtMallocCached.md→docs/03_api_ref/aclrtMallocCached.md+0-0
Rdocs/api_docs/aclrtMallocConfig.md→docs/03_api_ref/aclrtMallocConfig.md+0-0
Rdocs/api_docs/aclrtMallocForTaskScheduler.md→docs/03_api_ref/aclrtMallocForTaskScheduler.md+0-0
Rdocs/api_docs/aclrtMallocHost.md→docs/03_api_ref/aclrtMallocHost.md+0-0
Rdocs/api_docs/aclrtMallocHostWithCfg.md→docs/03_api_ref/aclrtMallocHostWithCfg.md+0-0
Rdocs/api_docs/aclrtMallocPhysical.md→docs/03_api_ref/aclrtMallocPhysical.md+0-0
Rdocs/api_docs/aclrtMallocWithCfg.md→docs/03_api_ref/aclrtMallocWithCfg.md+0-0
Rdocs/api_docs/aclrtMapMem.md→docs/03_api_ref/aclrtMapMem.md+0-0
Rdocs/api_docs/aclrtMbuf.md→docs/03_api_ref/aclrtMbuf.md+0-0
Rdocs/api_docs/aclrtMemAccessDesc.md→docs/03_api_ref/aclrtMemAccessDesc.md+0-0
Rdocs/api_docs/aclrtMemAccessFlags.md→docs/03_api_ref/aclrtMemAccessFlags.md+0-0
Rdocs/api_docs/aclrtMemAllocManaged.md→docs/03_api_ref/aclrtMemAllocManaged.md+0-0
Rdocs/api_docs/aclrtMemAllocationType.md→docs/03_api_ref/aclrtMemAllocationType.md+0-0
Rdocs/api_docs/aclrtMemAttr.md→docs/03_api_ref/aclrtMemAttr.md+0-0
Rdocs/api_docs/aclrtMemExportToShareableHandle.md→docs/03_api_ref/aclrtMemExportToShareableHandle.md+0-0
Rdocs/api_docs/aclrtMemExportToShareableHandleV2.md→docs/03_api_ref/aclrtMemExportToShareableHandleV2.md+0-0
Rdocs/api_docs/aclrtMemFlush.md→docs/03_api_ref/aclrtMemFlush.md+0-0
Rdocs/api_docs/aclrtMemGetAccess.md→docs/03_api_ref/aclrtMemGetAccess.md+0-0
Rdocs/api_docs/aclrtMemGetAddressRange.md→docs/03_api_ref/aclrtMemGetAddressRange.md+0-0
Rdocs/api_docs/aclrtMemGetAllocationGranularity.md→docs/03_api_ref/aclrtMemGetAllocationGranularity.md+0-0
Rdocs/api_docs/aclrtMemGetAllocationPropertiesFromHandle.md→docs/03_api_ref/aclrtMemGetAllocationPropertiesFromHandle.md+0-0
Rdocs/api_docs/aclrtMemGranularityOptions.md→docs/03_api_ref/aclrtMemGranularityOptions.md+0-0
Rdocs/api_docs/aclrtMemHandleType.md→docs/03_api_ref/aclrtMemHandleType.md+0-0
Rdocs/api_docs/aclrtMemImportFromShareableHandle.md→docs/03_api_ref/aclrtMemImportFromShareableHandle.md+0-0
Rdocs/api_docs/aclrtMemImportFromShareableHandleV2.md→docs/03_api_ref/aclrtMemImportFromShareableHandleV2.md+0-0
Rdocs/api_docs/aclrtMemInvalidate.md→docs/03_api_ref/aclrtMemInvalidate.md+0-0
Rdocs/api_docs/aclrtMemLocation.md→docs/03_api_ref/aclrtMemLocation.md+0-0
Rdocs/api_docs/aclrtMemLocationType.md→docs/03_api_ref/aclrtMemLocationType.md+0-0
Rdocs/api_docs/aclrtMemMallocPolicy.md→docs/03_api_ref/aclrtMemMallocPolicy.md+0-0
Rdocs/api_docs/aclrtMemManagedAdvise.md→docs/03_api_ref/aclrtMemManagedAdvise.md+35-35
Rdocs/api_docs/aclrtMemManagedAdviseType.md→docs/03_api_ref/aclrtMemManagedAdviseType.md+57-57
Rdocs/api_docs/aclrtMemManagedGetAttr.md→docs/03_api_ref/aclrtMemManagedGetAttr.md+39-39
Rdocs/api_docs/aclrtMemManagedGetAttrs.md→docs/03_api_ref/aclrtMemManagedGetAttrs.md+40-40
Rdocs/api_docs/aclrtMemManagedLocation.md→docs/03_api_ref/aclrtMemManagedLocation.md+11-11
Rdocs/api_docs/aclrtMemManagedLocationType.md→docs/03_api_ref/aclrtMemManagedLocationType.md+23-23
Rdocs/api_docs/aclrtMemManagedRangeAttribute.md→docs/03_api_ref/aclrtMemManagedRangeAttribute.md+93-93
Rdocs/api_docs/aclrtMemP2PMap.md→docs/03_api_ref/aclrtMemP2PMap.md+0-0
Rdocs/api_docs/aclrtMemPool.md→docs/03_api_ref/aclrtMemPool.md+0-0
Rdocs/api_docs/aclrtMemPoolAttr.md→docs/03_api_ref/aclrtMemPoolAttr.md+0-0
Rdocs/api_docs/aclrtMemPoolCreate.md→docs/03_api_ref/aclrtMemPoolCreate.md+0-0
Rdocs/api_docs/aclrtMemPoolDestroy.md→docs/03_api_ref/aclrtMemPoolDestroy.md+0-0
Rdocs/api_docs/aclrtMemPoolFreeAsync.md→docs/03_api_ref/aclrtMemPoolFreeAsync.md+29-29
Rdocs/api_docs/aclrtMemPoolGetAttr.md→docs/03_api_ref/aclrtMemPoolGetAttr.md+0-0
Rdocs/api_docs/aclrtMemPoolMallocAsync.md→docs/03_api_ref/aclrtMemPoolMallocAsync.md+32-32
Rdocs/api_docs/aclrtMemPoolProps.md→docs/03_api_ref/aclrtMemPoolProps.md+0-0
Rdocs/api_docs/aclrtMemPoolSetAttr.md→docs/03_api_ref/aclrtMemPoolSetAttr.md+0-0
Rdocs/api_docs/aclrtMemRetainAllocationHandle.md→docs/03_api_ref/aclrtMemRetainAllocationHandle.md+0-0
Rdocs/api_docs/aclrtMemSetAccess.md→docs/03_api_ref/aclrtMemSetAccess.md+0-0
Rdocs/api_docs/aclrtMemSetPidToShareableHandle.md→docs/03_api_ref/aclrtMemSetPidToShareableHandle.md+0-0
Rdocs/api_docs/aclrtMemSetPidToShareableHandleV2.md→docs/03_api_ref/aclrtMemSetPidToShareableHandleV2.md+0-0
Rdocs/api_docs/aclrtMemSharedHandleType.md→docs/03_api_ref/aclrtMemSharedHandleType.md+0-0
Rdocs/api_docs/aclrtMemUceInfo.md→docs/03_api_ref/aclrtMemUceInfo.md+0-0
Rdocs/api_docs/aclrtMemUceRepair.md→docs/03_api_ref/aclrtMemUceRepair.md+0-0
Rdocs/api_docs/aclrtMemUsageInfo.md→docs/03_api_ref/aclrtMemUsageInfo.md+0-0
Rdocs/api_docs/aclrtMemcpy.md→docs/03_api_ref/aclrtMemcpy.md+0-0
Rdocs/api_docs/aclrtMemcpy2d.md→docs/03_api_ref/aclrtMemcpy2d.md+0-0
Rdocs/api_docs/aclrtMemcpy2dAsync.md→docs/03_api_ref/aclrtMemcpy2dAsync.md+0-0
Rdocs/api_docs/aclrtMemcpyAsync.md→docs/03_api_ref/aclrtMemcpyAsync.md+0-0
Rdocs/api_docs/aclrtMemcpyAsyncWithCondition.md→docs/03_api_ref/aclrtMemcpyAsyncWithCondition.md+0-0
Rdocs/api_docs/aclrtMemcpyAsyncWithDesc.md→docs/03_api_ref/aclrtMemcpyAsyncWithDesc.md+0-0
Rdocs/api_docs/aclrtMemcpyAsyncWithOffset.md→docs/03_api_ref/aclrtMemcpyAsyncWithOffset.md+0-0
Rdocs/api_docs/aclrtMemcpyBatch.md→docs/03_api_ref/aclrtMemcpyBatch.md+0-0
Rdocs/api_docs/aclrtMemcpyBatchAsync.md→docs/03_api_ref/aclrtMemcpyBatchAsync.md+0-0
Rdocs/api_docs/aclrtMemcpyBatchAttr.md→docs/03_api_ref/aclrtMemcpyBatchAttr.md+0-0
Rdocs/api_docs/aclrtMemcpyKind.md→docs/03_api_ref/aclrtMemcpyKind.md+0-0
Rdocs/api_docs/aclrtMemset.md→docs/03_api_ref/aclrtMemset.md+0-0
Rdocs/api_docs/aclrtMemsetAsync.md→docs/03_api_ref/aclrtMemsetAsync.md+0-0
Rdocs/api_docs/aclrtNormalDisInfo.md→docs/03_api_ref/aclrtNormalDisInfo.md+0-0
Rdocs/api_docs/aclrtNotify.md→docs/03_api_ref/aclrtNotify.md+0-0
Rdocs/api_docs/aclrtNotifyBatchReset.md→docs/03_api_ref/aclrtNotifyBatchReset.md+0-0
Rdocs/api_docs/aclrtNotifyGetExportKey.md→docs/03_api_ref/aclrtNotifyGetExportKey.md+0-0
Rdocs/api_docs/aclrtNotifyImportByKey.md→docs/03_api_ref/aclrtNotifyImportByKey.md+0-0
Rdocs/api_docs/aclrtNotifySetImportPid.md→docs/03_api_ref/aclrtNotifySetImportPid.md+0-0
Rdocs/api_docs/aclrtNotifySetImportPidInterServer.md→docs/03_api_ref/aclrtNotifySetImportPidInterServer.md+0-0
Rdocs/api_docs/aclrtParamHandle.md→docs/03_api_ref/aclrtParamHandle.md+0-0
Rdocs/api_docs/aclrtPeekAtLastError.md→docs/03_api_ref/aclrtPeekAtLastError.md+0-0
Rdocs/api_docs/aclrtPersistentTaskClean.md→docs/03_api_ref/aclrtPersistentTaskClean.md+0-0
Rdocs/api_docs/aclrtPhysicalMemProp.md→docs/03_api_ref/aclrtPhysicalMemProp.md+0-0
Rdocs/api_docs/aclrtPointerGetAttributes.md→docs/03_api_ref/aclrtPointerGetAttributes.md+0-0
Rdocs/api_docs/aclrtProcessHostFunc.md→docs/03_api_ref/aclrtProcessHostFunc.md+0-0
Rdocs/api_docs/aclrtProcessReport.md→docs/03_api_ref/aclrtProcessReport.md+0-0
Rdocs/api_docs/aclrtProfTrace.md→docs/03_api_ref/aclrtProfTrace.md+0-0
Rdocs/api_docs/aclrtPtrAttributes.md→docs/03_api_ref/aclrtPtrAttributes.md+0-0
Rdocs/api_docs/aclrtQueryDeviceStatus.md→docs/03_api_ref/aclrtQueryDeviceStatus.md+0-0
Rdocs/api_docs/aclrtQueryEventStatus.md→docs/03_api_ref/aclrtQueryEventStatus.md+0-0
Rdocs/api_docs/aclrtQueryEventWaitStatus.md→docs/03_api_ref/aclrtQueryEventWaitStatus.md+0-0
Rdocs/api_docs/aclrtQueryEvent（废弃）.md→docs/03_api_ref/aclrtQueryEvent（废弃）.md+0-0
Rdocs/api_docs/aclrtRandomNumAsync.md→docs/03_api_ref/aclrtRandomNumAsync.md+0-0
Rdocs/api_docs/aclrtRandomNumFuncParaInfo.md→docs/03_api_ref/aclrtRandomNumFuncParaInfo.md+0-0
Rdocs/api_docs/aclrtRandomNumFuncType.md→docs/03_api_ref/aclrtRandomNumFuncType.md+0-0
Rdocs/api_docs/aclrtRandomNumTaskInfo.md→docs/03_api_ref/aclrtRandomNumTaskInfo.md+0-0
Rdocs/api_docs/aclrtRandomParaInfo.md→docs/03_api_ref/aclrtRandomParaInfo.md+0-0
Rdocs/api_docs/aclrtRandomTaskUpdateAttr.md→docs/03_api_ref/aclrtRandomTaskUpdateAttr.md+0-0
Rdocs/api_docs/aclrtRecordEvent.md→docs/03_api_ref/aclrtRecordEvent.md+0-0
Rdocs/api_docs/aclrtRecordNotify.md→docs/03_api_ref/aclrtRecordNotify.md+0-0
Rdocs/api_docs/aclrtReduceAsync.md→docs/03_api_ref/aclrtReduceAsync.md+0-0
Rdocs/api_docs/aclrtReduceKind.md→docs/03_api_ref/aclrtReduceKind.md+0-0
Rdocs/api_docs/aclrtRegDeviceStateCallback.md→docs/03_api_ref/aclrtRegDeviceStateCallback.md+0-0
Rdocs/api_docs/aclrtRegStreamStateCallback.md→docs/03_api_ref/aclrtRegStreamStateCallback.md+0-0
Rdocs/api_docs/aclrtRegisterCpuFunc.md→docs/03_api_ref/aclrtRegisterCpuFunc.md+0-0
Rdocs/api_docs/aclrtReleaseMemAddress.md→docs/03_api_ref/aclrtReleaseMemAddress.md+0-0
Rdocs/api_docs/aclrtRepairError.md→docs/03_api_ref/aclrtRepairError.md+0-0
Rdocs/api_docs/aclrtReserveMemAddress.md→docs/03_api_ref/aclrtReserveMemAddress.md+0-0
Rdocs/api_docs/aclrtReserveMemAddressNoUCMemory.md→docs/03_api_ref/aclrtReserveMemAddressNoUCMemory.md+0-0
Rdocs/api_docs/aclrtResetDevice.md→docs/03_api_ref/aclrtResetDevice.md+0-0
Rdocs/api_docs/aclrtResetDeviceForce.md→docs/03_api_ref/aclrtResetDeviceForce.md+0-0
Rdocs/api_docs/aclrtResetDeviceResLimit.md→docs/03_api_ref/aclrtResetDeviceResLimit.md+0-0
Rdocs/api_docs/aclrtResetEvent.md→docs/03_api_ref/aclrtResetEvent.md+0-0
Rdocs/api_docs/aclrtResetOverflowStatus.md→docs/03_api_ref/aclrtResetOverflowStatus.md+0-0
Rdocs/api_docs/aclrtResetStreamResLimit.md→docs/03_api_ref/aclrtResetStreamResLimit.md+0-0
Rdocs/api_docs/aclrtRunMode.md→docs/03_api_ref/aclrtRunMode.md+0-0
Rdocs/api_docs/aclrtServerPid.md→docs/03_api_ref/aclrtServerPid.md+0-0
Rdocs/api_docs/aclrtSetBufDataLen.md→docs/03_api_ref/aclrtSetBufDataLen.md+0-0
Rdocs/api_docs/aclrtSetBufUserData.md→docs/03_api_ref/aclrtSetBufUserData.md+0-0
Rdocs/api_docs/aclrtSetCurrentContext.md→docs/03_api_ref/aclrtSetCurrentContext.md+0-0
Rdocs/api_docs/aclrtSetDevice.md→docs/03_api_ref/aclrtSetDevice.md+0-0
Rdocs/api_docs/aclrtSetDeviceResLimit.md→docs/03_api_ref/aclrtSetDeviceResLimit.md+0-0
Rdocs/api_docs/aclrtSetDeviceSatMode.md→docs/03_api_ref/aclrtSetDeviceSatMode.md+0-0
Rdocs/api_docs/aclrtSetDeviceTaskAbortCallback.md→docs/03_api_ref/aclrtSetDeviceTaskAbortCallback.md+0-0
Rdocs/api_docs/aclrtSetExceptionInfoCallback.md→docs/03_api_ref/aclrtSetExceptionInfoCallback.md+0-0
Rdocs/api_docs/aclrtSetGroup.md→docs/03_api_ref/aclrtSetGroup.md+0-0
Rdocs/api_docs/aclrtSetLabel.md→docs/03_api_ref/aclrtSetLabel.md+0-0
Rdocs/api_docs/aclrtSetMemcpyDesc.md→docs/03_api_ref/aclrtSetMemcpyDesc.md+0-0
Rdocs/api_docs/aclrtSetOpExecuteTimeOut.md→docs/03_api_ref/aclrtSetOpExecuteTimeOut.md+0-0
Rdocs/api_docs/aclrtSetOpExecuteTimeOutV2.md→docs/03_api_ref/aclrtSetOpExecuteTimeOutV2.md+0-0
Rdocs/api_docs/aclrtSetOpExecuteTimeOutWithMs.md→docs/03_api_ref/aclrtSetOpExecuteTimeOutWithMs.md+0-0
Rdocs/api_docs/aclrtSetOpWaitTimeout.md→docs/03_api_ref/aclrtSetOpWaitTimeout.md+0-0
Rdocs/api_docs/aclrtSetStreamAttribute.md→docs/03_api_ref/aclrtSetStreamAttribute.md+0-0
Rdocs/api_docs/aclrtSetStreamConfigOpt.md→docs/03_api_ref/aclrtSetStreamConfigOpt.md+0-0
Rdocs/api_docs/aclrtSetStreamFailureMode.md→docs/03_api_ref/aclrtSetStreamFailureMode.md+0-0
Rdocs/api_docs/aclrtSetStreamOverflowSwitch.md→docs/03_api_ref/aclrtSetStreamOverflowSwitch.md+0-0
Rdocs/api_docs/aclrtSetStreamResLimit.md→docs/03_api_ref/aclrtSetStreamResLimit.md+0-0
Rdocs/api_docs/aclrtSetSysParamOpt.md→docs/03_api_ref/aclrtSetSysParamOpt.md+0-0
Rdocs/api_docs/aclrtSetTsDevice.md→docs/03_api_ref/aclrtSetTsDevice.md+0-0
Rdocs/api_docs/aclrtSnapShotCallbackRegister.md→docs/03_api_ref/aclrtSnapShotCallbackRegister.md+0-0
Rdocs/api_docs/aclrtSnapShotCallbackUnregister.md→docs/03_api_ref/aclrtSnapShotCallbackUnregister.md+0-0
Rdocs/api_docs/aclrtSnapShotProcessBackup.md→docs/03_api_ref/aclrtSnapShotProcessBackup.md+0-0
Rdocs/api_docs/aclrtSnapShotProcessLock.md→docs/03_api_ref/aclrtSnapShotProcessLock.md+0-0
Rdocs/api_docs/aclrtSnapShotProcessRestore.md→docs/03_api_ref/aclrtSnapShotProcessRestore.md+0-0
Rdocs/api_docs/aclrtSnapShotProcessUnlock.md→docs/03_api_ref/aclrtSnapShotProcessUnlock.md+0-0
Rdocs/api_docs/aclrtSnapShotStage.md→docs/03_api_ref/aclrtSnapShotStage.md+0-0
Rdocs/api_docs/aclrtStream.md→docs/03_api_ref/aclrtStream.md+0-0
Rdocs/api_docs/aclrtStreamAbort.md→docs/03_api_ref/aclrtStreamAbort.md+0-0
Rdocs/api_docs/aclrtStreamAttr.md→docs/03_api_ref/aclrtStreamAttr.md+0-0
Rdocs/api_docs/aclrtStreamAttrValue.md→docs/03_api_ref/aclrtStreamAttrValue.md+0-0
Rdocs/api_docs/aclrtStreamConfigAttr.md→docs/03_api_ref/aclrtStreamConfigAttr.md+0-0
Rdocs/api_docs/aclrtStreamConfigHandle.md→docs/03_api_ref/aclrtStreamConfigHandle.md+0-0
Rdocs/api_docs/aclrtStreamGetFlags.md→docs/03_api_ref/aclrtStreamGetFlags.md+0-0
Rdocs/api_docs/aclrtStreamGetId.md→docs/03_api_ref/aclrtStreamGetId.md+0-0
Rdocs/api_docs/aclrtStreamGetPriority.md→docs/03_api_ref/aclrtStreamGetPriority.md+0-0
Rdocs/api_docs/aclrtStreamQuery.md→docs/03_api_ref/aclrtStreamQuery.md+0-0
Rdocs/api_docs/aclrtStreamStatus.md→docs/03_api_ref/aclrtStreamStatus.md+0-0
Rdocs/api_docs/aclrtStreamStop.md→docs/03_api_ref/aclrtStreamStop.md+0-0
Rdocs/api_docs/aclrtStreamWaitEvent.md→docs/03_api_ref/aclrtStreamWaitEvent.md+0-0
Rdocs/api_docs/aclrtStreamWaitEventWithTimeout.md→docs/03_api_ref/aclrtStreamWaitEventWithTimeout.md+0-0
Rdocs/api_docs/aclrtSubscribeHostFunc.md→docs/03_api_ref/aclrtSubscribeHostFunc.md+0-0
Rdocs/api_docs/aclrtSubscribeReport.md→docs/03_api_ref/aclrtSubscribeReport.md+0-0
Rdocs/api_docs/aclrtSwitchLabelByIndex.md→docs/03_api_ref/aclrtSwitchLabelByIndex.md+0-0
Rdocs/api_docs/aclrtSwitchStream.md→docs/03_api_ref/aclrtSwitchStream.md+0-0
Rdocs/api_docs/aclrtSynchronizeDevice.md→docs/03_api_ref/aclrtSynchronizeDevice.md+0-0
Rdocs/api_docs/aclrtSynchronizeDeviceWithTimeout.md→docs/03_api_ref/aclrtSynchronizeDeviceWithTimeout.md+0-0
Rdocs/api_docs/aclrtSynchronizeEvent.md→docs/03_api_ref/aclrtSynchronizeEvent.md+0-0
Rdocs/api_docs/aclrtSynchronizeEventWithTimeout.md→docs/03_api_ref/aclrtSynchronizeEventWithTimeout.md+0-0
Rdocs/api_docs/aclrtSynchronizeStream.md→docs/03_api_ref/aclrtSynchronizeStream.md+0-0
Rdocs/api_docs/aclrtSynchronizeStreamWithTimeout.md→docs/03_api_ref/aclrtSynchronizeStreamWithTimeout.md+0-0
Rdocs/api_docs/aclrtTaskGrp.md→docs/03_api_ref/aclrtTaskGrp.md+0-0
Rdocs/api_docs/aclrtTaskUpdateAsync.md→docs/03_api_ref/aclrtTaskUpdateAsync.md+0-0
Rdocs/api_docs/aclrtTaskUpdateInfo.md→docs/03_api_ref/aclrtTaskUpdateInfo.md+0-0
Rdocs/api_docs/aclrtTimeoutUs.md→docs/03_api_ref/aclrtTimeoutUs.md+0-0
Rdocs/api_docs/aclrtUnSubscribeHostFunc.md→docs/03_api_ref/aclrtUnSubscribeHostFunc.md+0-0
Rdocs/api_docs/aclrtUnSubscribeReport.md→docs/03_api_ref/aclrtUnSubscribeReport.md+0-0
Rdocs/api_docs/aclrtUniformDisInfo.md→docs/03_api_ref/aclrtUniformDisInfo.md+0-0
Rdocs/api_docs/aclrtUnmapMem.md→docs/03_api_ref/aclrtUnmapMem.md+0-0
Rdocs/api_docs/aclrtUnuseStreamResInCurrentThread.md→docs/03_api_ref/aclrtUnuseStreamResInCurrentThread.md+0-0
Rdocs/api_docs/aclrtUpdateTaskAttrId.md→docs/03_api_ref/aclrtUpdateTaskAttrId.md+0-0
Rdocs/api_docs/aclrtUpdateTaskAttrVal.md→docs/03_api_ref/aclrtUpdateTaskAttrVal.md+0-0
Rdocs/api_docs/aclrtUseStreamResInCurrentThread.md→docs/03_api_ref/aclrtUseStreamResInCurrentThread.md+0-0
Rdocs/api_docs/aclrtUtilizationInfo.md→docs/03_api_ref/aclrtUtilizationInfo.md+0-0
Rdocs/api_docs/aclrtUuid.md→docs/03_api_ref/aclrtUuid.md+0-0
Rdocs/api_docs/aclrtValueWait.md→docs/03_api_ref/aclrtValueWait.md+0-0
Rdocs/api_docs/aclrtValueWrite.md→docs/03_api_ref/aclrtValueWrite.md+0-0
Rdocs/api_docs/aclrtWaitAndResetNotify.md→docs/03_api_ref/aclrtWaitAndResetNotify.md+0-0
Rdocs/api_docs/aclsysGetCANNVersion.md→docs/03_api_ref/aclsysGetCANNVersion.md+0-0
Rdocs/api_docs/aclsysGetVersionNum.md→docs/03_api_ref/aclsysGetVersionNum.md+0-0
Rdocs/api_docs/aclsysGetVersionStr.md→docs/03_api_ref/aclsysGetVersionStr.md+0-0
Rdocs/api_docs/aclsysGetVesionNum.md→docs/03_api_ref/aclsysGetVesionNum.md+0-0
Rdocs/api_docs/aclsysGetVesionStr.md→docs/03_api_ref/aclsysGetVesionStr.md+0-0
Rdocs/api_docs/acltdtAddDataItem.md→docs/03_api_ref/acltdtAddDataItem.md+0-0
Rdocs/api_docs/acltdtAddQueueRoute.md→docs/03_api_ref/acltdtAddQueueRoute.md+0-0
Rdocs/api_docs/acltdtAllocBuf.md→docs/03_api_ref/acltdtAllocBuf.md+0-0
Rdocs/api_docs/acltdtAppendBufChain.md→docs/03_api_ref/acltdtAppendBufChain.md+0-0
Rdocs/api_docs/acltdtAttachQueue.md→docs/03_api_ref/acltdtAttachQueue.md+0-0
Rdocs/api_docs/acltdtBindQueueRoutes.md→docs/03_api_ref/acltdtBindQueueRoutes.md+0-0
Rdocs/api_docs/acltdtBuf.md→docs/03_api_ref/acltdtBuf.md+0-0
Rdocs/api_docs/acltdtCleanChannel.md→docs/03_api_ref/acltdtCleanChannel.md+0-0
Rdocs/api_docs/acltdtCopyBufRef.md→docs/03_api_ref/acltdtCopyBufRef.md+0-0
Rdocs/api_docs/acltdtCreateChannel.md→docs/03_api_ref/acltdtCreateChannel.md+0-0
Rdocs/api_docs/acltdtCreateChannelWithCapacity.md→docs/03_api_ref/acltdtCreateChannelWithCapacity.md+0-0
Rdocs/api_docs/acltdtCreateDataItem.md→docs/03_api_ref/acltdtCreateDataItem.md+0-0
Rdocs/api_docs/acltdtCreateDataset.md→docs/03_api_ref/acltdtCreateDataset.md+0-0
Rdocs/api_docs/acltdtCreateQueue.md→docs/03_api_ref/acltdtCreateQueue.md+0-0
Rdocs/api_docs/acltdtCreateQueueAttr.md→docs/03_api_ref/acltdtCreateQueueAttr.md+0-0
Rdocs/api_docs/acltdtCreateQueueRoute.md→docs/03_api_ref/acltdtCreateQueueRoute.md+0-0
Rdocs/api_docs/acltdtCreateQueueRouteList.md→docs/03_api_ref/acltdtCreateQueueRouteList.md+0-0
Rdocs/api_docs/acltdtCreateQueueRouteQueryInfo.md→docs/03_api_ref/acltdtCreateQueueRouteQueryInfo.md+0-0
Rdocs/api_docs/acltdtDataItem.md→docs/03_api_ref/acltdtDataItem.md+0-0
Rdocs/api_docs/acltdtDataset.md→docs/03_api_ref/acltdtDataset.md+0-0
Rdocs/api_docs/acltdtDequeue.md→docs/03_api_ref/acltdtDequeue.md+0-0
Rdocs/api_docs/acltdtDequeueData.md→docs/03_api_ref/acltdtDequeueData.md+0-0
Rdocs/api_docs/acltdtDestroyChannel.md→docs/03_api_ref/acltdtDestroyChannel.md+0-0
Rdocs/api_docs/acltdtDestroyDataItem.md→docs/03_api_ref/acltdtDestroyDataItem.md+0-0
Rdocs/api_docs/acltdtDestroyDataset.md→docs/03_api_ref/acltdtDestroyDataset.md+0-0
Rdocs/api_docs/acltdtDestroyQueue.md→docs/03_api_ref/acltdtDestroyQueue.md+0-0
Rdocs/api_docs/acltdtDestroyQueueAttr.md→docs/03_api_ref/acltdtDestroyQueueAttr.md+0-0
Rdocs/api_docs/acltdtDestroyQueueRoute.md→docs/03_api_ref/acltdtDestroyQueueRoute.md+0-0
Rdocs/api_docs/acltdtDestroyQueueRouteList.md→docs/03_api_ref/acltdtDestroyQueueRouteList.md+0-0
Rdocs/api_docs/acltdtDestroyQueueRouteQueryInfo.md→docs/03_api_ref/acltdtDestroyQueueRouteQueryInfo.md+0-0
Rdocs/api_docs/acltdtEnqueue.md→docs/03_api_ref/acltdtEnqueue.md+0-0
Rdocs/api_docs/acltdtEnqueueData.md→docs/03_api_ref/acltdtEnqueueData.md+0-0
Rdocs/api_docs/acltdtFreeBuf.md→docs/03_api_ref/acltdtFreeBuf.md+0-0
Rdocs/api_docs/acltdtGetBufChainNum.md→docs/03_api_ref/acltdtGetBufChainNum.md+0-0
Rdocs/api_docs/acltdtGetBufData.md→docs/03_api_ref/acltdtGetBufData.md+0-0
Rdocs/api_docs/acltdtGetBufDataLen.md→docs/03_api_ref/acltdtGetBufDataLen.md+0-0
Rdocs/api_docs/acltdtGetBufFromChain.md→docs/03_api_ref/acltdtGetBufFromChain.md+0-0
Rdocs/api_docs/acltdtGetBufUserData.md→docs/03_api_ref/acltdtGetBufUserData.md+0-0
Rdocs/api_docs/acltdtGetDataAddrFromItem.md→docs/03_api_ref/acltdtGetDataAddrFromItem.md+0-0
Rdocs/api_docs/acltdtGetDataItem.md→docs/03_api_ref/acltdtGetDataItem.md+0-0
Rdocs/api_docs/acltdtGetDataSizeFromItem.md→docs/03_api_ref/acltdtGetDataSizeFromItem.md+0-0
Rdocs/api_docs/acltdtGetDataTypeFromItem.md→docs/03_api_ref/acltdtGetDataTypeFromItem.md+0-0
Rdocs/api_docs/acltdtGetDatasetName.md→docs/03_api_ref/acltdtGetDatasetName.md+0-0
Rdocs/api_docs/acltdtGetDatasetSize.md→docs/03_api_ref/acltdtGetDatasetSize.md+0-0
Rdocs/api_docs/acltdtGetDimNumFromItem.md→docs/03_api_ref/acltdtGetDimNumFromItem.md+0-0
Rdocs/api_docs/acltdtGetDimsFromItem.md→docs/03_api_ref/acltdtGetDimsFromItem.md+0-0
Rdocs/api_docs/acltdtGetQueueAttr.md→docs/03_api_ref/acltdtGetQueueAttr.md+0-0
Rdocs/api_docs/acltdtGetQueueRoute.md→docs/03_api_ref/acltdtGetQueueRoute.md+0-0
Rdocs/api_docs/acltdtGetQueueRouteNum.md→docs/03_api_ref/acltdtGetQueueRouteNum.md+0-0
Rdocs/api_docs/acltdtGetQueueRouteParam.md→docs/03_api_ref/acltdtGetQueueRouteParam.md+0-0
Rdocs/api_docs/acltdtGetSliceInfoFromItem.md→docs/03_api_ref/acltdtGetSliceInfoFromItem.md+0-0
Rdocs/api_docs/acltdtGetTensorTypeFromItem.md→docs/03_api_ref/acltdtGetTensorTypeFromItem.md+0-0
Rdocs/api_docs/acltdtGrantQueue.md→docs/03_api_ref/acltdtGrantQueue.md+0-0
Rdocs/api_docs/acltdtQueryChannelSize.md→docs/03_api_ref/acltdtQueryChannelSize.md+0-0
Rdocs/api_docs/acltdtQueryQueueRoutes.md→docs/03_api_ref/acltdtQueryQueueRoutes.md+0-0
Rdocs/api_docs/acltdtQueueAttr.md→docs/03_api_ref/acltdtQueueAttr.md+0-0
Rdocs/api_docs/acltdtQueueAttrType.md→docs/03_api_ref/acltdtQueueAttrType.md+0-0
Rdocs/api_docs/acltdtQueueRoute.md→docs/03_api_ref/acltdtQueueRoute.md+0-0
Rdocs/api_docs/acltdtQueueRouteList.md→docs/03_api_ref/acltdtQueueRouteList.md+0-0
Rdocs/api_docs/acltdtQueueRouteParamType.md→docs/03_api_ref/acltdtQueueRouteParamType.md+0-0
Rdocs/api_docs/acltdtQueueRouteQueryInfo.md→docs/03_api_ref/acltdtQueueRouteQueryInfo.md+0-0
Rdocs/api_docs/acltdtQueueRouteQueryInfoParamType.md→docs/03_api_ref/acltdtQueueRouteQueryInfoParamType.md+0-0
Rdocs/api_docs/acltdtQueueRouteQueryMode.md→docs/03_api_ref/acltdtQueueRouteQueryMode.md+0-0
Rdocs/api_docs/acltdtReceiveTensor.md→docs/03_api_ref/acltdtReceiveTensor.md+0-0
Rdocs/api_docs/acltdtSendTensor.md→docs/03_api_ref/acltdtSendTensor.md+0-0
Rdocs/api_docs/acltdtSetBufDataLen.md→docs/03_api_ref/acltdtSetBufDataLen.md+0-0
Rdocs/api_docs/acltdtSetBufUserData.md→docs/03_api_ref/acltdtSetBufUserData.md+0-0
Rdocs/api_docs/acltdtSetQueueAttr.md→docs/03_api_ref/acltdtSetQueueAttr.md+0-0
Rdocs/api_docs/acltdtSetQueueRouteQueryInfo.md→docs/03_api_ref/acltdtSetQueueRouteQueryInfo.md+0-0
Rdocs/api_docs/acltdtStopChannel.md→docs/03_api_ref/acltdtStopChannel.md+0-0
Rdocs/api_docs/acltdtTensorType.md→docs/03_api_ref/acltdtTensorType.md+0-0
Rdocs/api_docs/acltdtUnbindQueueRoutes.md→docs/03_api_ref/acltdtUnbindQueueRoutes.md+0-0
Adocs/03_api_ref/api_ref.md+634-0
Rdocs/api_docs/figures/Device-Context-Stream之间的关系.png→docs/03_api_ref/figures/Device-Context-Stream之间的关系.png+0-0
Rdocs/api_docs/figures/SOMA_Reuse.png→docs/03_api_ref/figures/SOMA_Reuse.png+0-0
Rdocs/api_docs/figures/aclmdlDataset类型与aclDataBuffer类型的关系-2.png→docs/03_api_ref/figures/aclmdlDataset类型与aclDataBuffer类型的关系-2.png+0-0
Rdocs/api_docs/figures/aclmdlDataset类型与aclDataBuffer类型的关系.png→docs/03_api_ref/figures/aclmdlDataset类型与aclDataBuffer类型的关系.png+0-0
Rdocs/api_docs/figures/zh-cn_image_0000001264922118.png→docs/03_api_ref/figures/zh-cn_image_0000001264922118.png+0-0
Rdocs/api_docs/figures/zh-cn_image_0000002320610726.png→docs/03_api_ref/figures/zh-cn_image_0000002320610726.png+0-0
Rdocs/api_docs/figures/主要接口调用流程.png→docs/03_api_ref/figures/主要接口调用流程.png+0-0
Rdocs/api_docs/figures/同步等待流程_多Device场景.png→docs/03_api_ref/figures/同步等待流程_多Device场景.png+0-0
Rdocs/api_docs/figures/接口调用流程图-0.png→docs/03_api_ref/figures/接口调用流程图-0.png+0-0
Rdocs/api_docs/figures/接口调用流程图-1.png→docs/03_api_ref/figures/接口调用流程图-1.png+0-0
Rdocs/api_docs/figures/接口调用流程图.png→docs/03_api_ref/figures/接口调用流程图.png+0-0
Rdocs/api_docs/figures/逻辑架构图-3.png→docs/03_api_ref/figures/逻辑架构图-3.png+0-0
Rdocs/api_docs/figures/逻辑架构图-4.png→docs/03_api_ref/figures/逻辑架构图-4.png+0-0
Rdocs/api_docs/figures/逻辑架构图.png→docs/03_api_ref/figures/逻辑架构图.png+0-0
Rdocs/api_docs/msproftx扩展接口.md→docs/03_api_ref/msproftx扩展接口.md+0-0
Rdocs/api_docs/使用说明.md→docs/03_api_ref/使用说明.md+0-0
Rdocs/api_docs/共享Buffer管理.md→docs/03_api_ref/共享Buffer管理.md+0-0
Rdocs/api_docs/共享Buffer管理（预留-暂不支持）.md→docs/03_api_ref/共享Buffer管理（预留-暂不支持）.md+0-0
Rdocs/api_docs/共享队列管理.md→docs/03_api_ref/共享队列管理.md+0-0
Rdocs/api_docs/其他接口.md→docs/03_api_ref/其他接口.md+0-0
Rdocs/api_docs/内存使用说明.md→docs/03_api_ref/内存使用说明.md+0-0
Rdocs/api_docs/内存管理.md→docs/03_api_ref/内存管理.md+0-0
Rdocs/api_docs/初始化-去初始化.md→docs/03_api_ref/初始化-去初始化.md+0-0
Rdocs/api_docs/同步-异步API说明.md→docs/03_api_ref/同步-异步API说明.md+0-0
Rdocs/api_docs/头文件-库文件说明.md→docs/03_api_ref/头文件-库文件说明.md+0-0
Rdocs/api_docs/废弃接口-返回码列表.md→docs/03_api_ref/废弃接口-返回码列表.md+0-0
Rdocs/api_docs/异常处理.md→docs/03_api_ref/异常处理.md+0-0
Rdocs/api_docs/快照管理.md→docs/03_api_ref/快照管理.md+0-0
Rdocs/api_docs/执行控制.md→docs/03_api_ref/执行控制.md+0-0
Rdocs/api_docs/扩展接口使用说明.md→docs/03_api_ref/扩展接口使用说明.md+0-0
Rdocs/api_docs/数据传输.md→docs/03_api_ref/数据传输.md+0-0
Rdocs/api_docs/数据类型及其操作接口.md→docs/03_api_ref/数据类型及其操作接口.md+0-0
Rdocs/api_docs/数据采集说明.md→docs/03_api_ref/数据采集说明.md+0-0
Rdocs/api_docs/概念及使用说明.md→docs/03_api_ref/概念及使用说明.md+0-0
Rdocs/api_docs/模型运行实例管理.md→docs/03_api_ref/模型运行实例管理.md+0-0
Rdocs/api_docs/简介.md→docs/03_api_ref/简介.md+0-0
Rdocs/api_docs/算力Group查询与设置.md→docs/03_api_ref/算力Group查询与设置.md+0-0
Rdocs/api_docs/订阅接口使用说明.md→docs/03_api_ref/订阅接口使用说明.md+0-0
Rdocs/api_docs/订阅算子信息.md→docs/03_api_ref/订阅算子信息.md+0-0
Rdocs/api_docs/运行时配置.md→docs/03_api_ref/运行时配置.md+0-0