Update the programming guide and merged the .md files.-runtime-AtomGit

已合并

Update the programming guide and merged the .md files. #1561

ycm0028创建于 4月14日从已删除 :master 合入到 cann/runtime： master

Update the programming guide and merged the .md files. #1561

已合并

ycm0028创建于 4月14日

从已删除 :master合入到cann/runtime master

收起

docs
- 02_dev_guide
  - 00_dev_guide.md
  - 01_初始化.md
  - 02_内存管理.md
  - 03-01_异步任务总述.md
  - 03-02_Stream管理.md
  - 03-03_Kernel加载与执行.md
  - 03-04_系统任务.md
  - 03-05_Event管理.md
  - 03-06_Notify管理.md
  - 03-07_内存语义同步.md
  - 03_异步任务执行.md
  - 04-01_单流捕获.md
  - 04-02_跨流捕获.md
  - 04-03_任务更新.md
  - 04_ACL-Graph.md
  - 05_多设备编程.md
  - 06_进程间通信.md
  - 07_运行时核资源控制.md
  - 08_配置AI-Core栈空间大小.md
  - 09_异常处理.md
  - Device内存使用.md
  - Event同步.md
  - Event查询.md
  - Event概念.md
  - Event的创建与销毁.md
  - Event等待.md
  - Event管理.md
  - Host回调任务.md
  - Host锁页内存使用.md
  - Persistent流.md
  - Stream创建与销毁.md
  - Stream和Event的行为.md
  - Stream概念.md
  - Stream管理.md
  - dev_guide.md
  - 内存管理.md
  - 内存管理总述.md
  - 多Device选择.md
  - 多设备编程.md
  - 多设备编程总述.md
  - 异常处理.md
  - 异步任务执行.md
  - 显式同步.md
  - 记录Event时间戳.md
  - 配置Stream优先级.md
  - 配置任务遇错即停.md
  - 默认Stream.md
- README.md

共 48 个文件变更+752-780

Adocs/02_dev_guide/00_dev_guide.md+22-0

@@ -0,0 +1,22 @@
	1	+# Runtime编程指南
	2	+
	3	+- ## [1. 初始化](01_初始化.md)
	4	+- ## [2. 内存管理](02_内存管理.md)
	5	+- ## [3. 异步任务执行](03_异步任务执行.md)
	6	+ - ### [3.1 异步任务总述](03-01_异步任务总述.md)
	7	+ - ### [3.2 Stream管理](03-02_Stream管理.md)
	8	+ - ### [3.3 Kernel加载与执行](03-03_Kernel加载与执行.md)
	9	+ - ### [3.4 系统任务](03-04_系统任务.md)
	10	+ - ### [3.5 Event管理](03-05_Event管理.md)
	11	+ - ### [3.6 Notify管理](03-06_Notify管理.md)
	12	+ - ### [3.7 内存语义同步](03-07_内存语义同步.md)
	13	+- ## [4. ACL Graph](04_ACL-Graph.md)
	14	+ - ### [4.1 单流捕获](04-01_单流捕获.md)
	15	+ - ### [4.2 跨流捕获](04-02_跨流捕获.md)
	16	+ - ### [4.3 任务更新](04-03_任务更新.md)
	17	+- ## [5. 多设备编程](05_多设备编程.md)
	18	+- ## [6. 进程间通信](06_进程间通信.md)
	19	+- ## [7. 运行时核资源控制](07_运行时核资源控制.md)
	20	+- ## [8. 配置AI Core栈空间大小](08_配置AI-Core栈空间大小.md)
	21	+- ## [9. 异常处理](09_异常处理.md)
	22	+

Rdocs/02_dev_guide/初始化.md→docs/02_dev_guide/01_初始化.md+0-0

文件重命名但无更改。

Rdocs/02_dev_guide/虚拟内存管理.md→docs/02_dev_guide/02_内存管理.md+154-1

	@@ -1,4 +1,157 @@
1	-# 虚拟内存管理	1	+# 内存管理
		2	+
		3	+## 内存管理总述
		4	+
		5	+在昇腾异构计算架构中，系统由主机（Host）和设备（Device）组成。Host和Device各自拥有独立的内存，Host内存是指AI处理器所在服务器的主机内存（即CPU内存），而Device内存则是指AI处理器自带的设备内存。
		6	+
		7	+内存管理中要做好的两件事是：
		8	+
		9	+1. 可以访问内存：Runtime提供了一套内存管理API，使开发者能够高效便捷地编写应用程序中的内存管理代码。由于Host和Device的内存相互独立，Runtime提供了专门的接口来分别申请和释放Host内存及Device内存。例如，申请和释放Host内存的接口为aclrtMallocHost和aclrtFreeHost，而申请和释放Device内存的接口为aclrtMalloc和aclrtFree。
		10	+2. 高效访问内存：为了实现最佳的内存访问性能，需要将数据存储在对应的内存中，例如算子在Device上执行过程中，访问Device上的数据性能要远高于访问Host侧的。为此，Runtime提供了Host与Device之间互相拷贝内存的接口，支持同步和异步方式，例如aclrtMemcpy和aclrtMemcpyAsync等，以便开发者更好地规划数据的存储与访问。
		11	+
		12	+## Device内存使用
		13	+
		14	+在昇腾异构计算编程中，典型的使用场景是：通过aclrtMallocHost接口申请Host内存，通过aclrtMalloc接口申请Device内存，通过aclrtMemcpy（同步）/aclrtMemcpyAsync（异步）接口将数据从Host拷贝到Device上，算子执行过程中使用Device内存进行计算并保存结果。
		15	+
		16	+以下是一段简单的示例代码。在示例代码中，两个张量从Host内存被拷贝到Device内存，在Device侧完成计算，再将结果从Device内存拷贝到Host内存：
		17	+
		18	+```
		19	+int main(void)
		20	+{
		21	+ int32_t deviceId = 0;
		22	+ int64_t N = 16;
		23	+ const size_t bytes = static_cast<size_t>(N) * sizeof(float);
		24	+ aclrtStream stream = nullptr;
		25	+
		26	+ // STEP 1: 初始化、Stream创建
		27	+ aclInit(nullptr);
		28	+ aclrtSetDevice(deviceId);
		29	+ aclrtCreateStream(&stream);
		30	+
		31	+ // STEP 2: 申请Host内存
		32	+ void* hostA = nullptr;
		33	+ void* hostB = nullptr;
		34	+ void* hostOut = nullptr;
		35	+
		36	+ aclrtMallocHost(&hostA, bytes);
		37	+ aclrtMallocHost(&hostB, bytes);
		38	+ aclrtMallocHost(&hostOut, bytes);
		39	+
		40	+ // 输入数据初始化
		41	+ ...
		42	+
		43	+ // STEP 3: 申请Device内存
		44	+ void* deviceA = nullptr;
		45	+ void* deviceB = nullptr;
		46	+ void* deviceOut = nullptr;
		47	+
		48	+ // 第三个参数aclrtMemMallocPolicy表示申请内存时的内存页分配策略
		49	+ // ACL_MEM_MALLOC_HUGE_FIRST 表示大页内存优先，其余定义参考API文档
		50	+ aclrtMalloc(&deviceA, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
		51	+ aclrtMalloc(&deviceB, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
		52	+ aclrtMalloc(&deviceOut, bytes, ACL_MEM_MALLOC_HUGE_FIRST);
		53	+
		54	+ // STEP 4: 将输入数据从Host内存传输到Device内存
		55	+ aclrtMemcpy(deviceA, bytes, hostA, bytes, ACL_MEMCPY_HOST_TO_DEVICE);
		56	+ aclrtMemcpy(deviceB, bytes, hostB, bytes, ACL_MEMCPY_HOST_TO_DEVICE);
		57	+
		58	+ // STEP 5: 执行计算操作，比如aclnnAdd，将计算结果保存在Device内存
		59	+ // ...
		60	+
		61	+ // STEP 6: 将计算结果从Device内存传输到Host内存
		62	+ aclrtMemcpy(hostOut, bytes, deviceOut, bytes, ACL_MEMCPY_DEVICE_TO_HOST);
		63	+
		64	+ // STEP 7: 在Host侧处理计算结果
		65	+ // ...
		66	+
		67	+ // STEP 8: 资源清理
		68	+ // STEP 8.1: 释放Host和Device内存
		69	+ if (deviceA) (void)aclrtFree(deviceA);
		70	+ if (deviceB) (void)aclrtFree(deviceB);
		71	+ if (deviceOut) (void)aclrtFree(deviceOut);
		72	+ if (hostA) (void)aclrtFreeHost(hostA);
		73	+ if (hostB) (void)aclrtFreeHost(hostB);
		74	+ if (hostOut) (void)aclrtFreeHost(hostOut);
		75	+
		76	+ // STEP 8.2: 清理设备和进程资源
		77	+ if (stream) (void)aclrtDestroyStream(stream);
		78	+ (void)aclrtResetDevice(deviceId);
		79	+ (void)aclFinalize();
		80	+
		81	+ return 0;
		82	+}
		83	+```
		84	+
		85	+## Host锁页内存使用
		86	+
		87	+在CANN编程框架中，Host内存可以是Pageable内存，也可以是Page-Locked内存（也称锁页内存或Pinned内存）：
		88	+
		89	+- Pageable内存，由操作系统统一管理。开发者可使用malloc、mmap等传统接口申请内存，使用free、munmap等接口释放内存。当内存压力较大时，Pageable内存会被换出到后备存储提供的交换空间中。当Pageable内存中的数据被传输到Device时，数据首先会被复制到缓冲区，然后通过DMA（Direct Memory Access）通道传输到Device。
		90	+- Page-Locked内存，即锁页内存。开发者需要用Runtime提供的API进行锁页内存的申请和释放，例如aclrtMallocHost、aclrtFreeHost等接口。对于锁页内存，虚拟页与物理页的映射关系固定，在其生命周期内不会被换出至交换空间。当锁页内存中的数据被传输至Device时，直接通过DMA通道传输，无需经过缓冲区，传输性能更优。
		91	+
		92	+ 在Runtime中，使用锁页内存的好处如下：
		93	+
		94	+ - 设备可以直接通过DMA访问主机内存，无需经过缓冲区，可以提供更好的传输性能；
		95	+ - 数据搬运过程无需CPU参与，可以实现数据的异步传输；
		96	+ - 数据的异步传输，使传输过程和计算过程可以相互掩盖，减少传输+计算的整体时长。
		97	+
		98	+锁页内存可直接通过aclrtMallocHost接口申请，示例代码如下。如果需要在申请时指定内存的其他配置，例如自定义模块ID、配置VA（virtual address）一致性等，也可以使用aclrtMallocHostWithCfg接口申请。
		99	+
		100	+```
		101	+// 资源初始化
		102	+......
		103	+
		104	+// 申请锁页内存
		105	+void *hostPtr = NULL;
		106	+aclrtMallocHost(&hostPtr, size);
		107	+
		108	+// 申请Device内存
		109	+void *devicePtr = NULL;
		110	+aclrtMalloc(&devicePtr, size, ACL_MEM_MALLOC_NORMAL_ONLY);
		111	+
		112	+// 内存初始化
		113	+......
		114	+
		115	+// 异步H2D
		116	+aclrtMemcpyAsync(devicePtr, size, hostPtr, size, ACL_MEMCPY_HOST_TO_DEVICE, stream);
		117	+
		118	+// 等待异步拷贝完成
		119	+aclrtSynchronizeStream(stream);
		120	+
		121	+// 资源释放
		122	+if (devicePtr) (void)aclrtFree(devicePtr);
		123	+if (hostPtr) (void)aclrtFreeHost(hostPtr);
		124	+......
		125	+```
		126	+
		127	+若使用malloc/mmap等内存管理接口申请Pageable内存，当前Runtime提供了aclrtHostRegisterV2接口，用于将Pageable内存转换为锁页内存，供Device访问。请注意，当OS内核版本为5.10或更低时，此方法会导致异常，此时应通过aclrtMallocHost接口申请锁页内存。
		128	+
		129	+```
		130	+// 申请Pageable内存
		131	+void *hostPtr = malloc(size);
		132	+
		133	+// 注册为锁页内存(ACL_HOST_REG_PINNED)，并映射到Device(ACL_HOST_REG_MAPPED)
		134	+aclrtHostRegisterV2(hostPtr, size, ACL_HOST_REG_PINNED\|ACL_HOST_REG_MAPPED);
		135	+
		136	+// 获取Device地址
		137	+void *devicePtr = NULL;
		138	+aclrtHostGetDevicePointer(hostPtr, &devicePtr, 0);
		139	+
		140	+// 任务通过Device地址访问对应内存
		141	+...
		142	+
		143	+// 资源释放
		144	+aclrtHostUnregister(hostPtr);
		145	+free(hostPtr);
		146	+```
		147	+
		148	+若涉及Host内存的VA（virtual address）一致性：
		149	+
		150	+- Host内存可以通过aclrtHostRegister接口注册到Device，根据Host指针映射得到Device指针，供Device使用；默认情况下，对于同一块Host内存，Host和Device看到的虚拟地址是不同的；
		151	+- 如果需要Host和Device的虚拟地址保持一致，可以使用aclrtMallocHostWithCfg接口申请锁页内存，并在aclrtMallocConfig参数中指定ACL\_RT\_MEM\_ATTR\_VA\_FLAG属性，后续再注册获取到的Device虚拟地址即可与Host虚拟地址一致。
		152	+
		153	+
		154	+## 虚拟内存管理
2		155
3	Runtime提供了一套虚拟内存管理的API接口，可供开发者更加精细化的管理其内存，例如：可以提前申请好大片的连续的虚拟内存，以简化后续的管理使用，同时可以按需在使用过程中申请物理内存并与虚拟内存做地址映射，即可以精细化管理，也可以最大化利用物理内存；此外还可以在进程之间利用虚拟内存管理中的handle来实现物理内存共享。	156	Runtime提供了一套虚拟内存管理的API接口，可供开发者更加精细化的管理其内存，例如：可以提前申请好大片的连续的虚拟内存，以简化后续的管理使用，同时可以按需在使用过程中申请物理内存并与虚拟内存做地址映射，即可以精细化管理，也可以最大化利用物理内存；此外还可以在进程之间利用虚拟内存管理中的handle来实现物理内存共享。
4		157

Rdocs/02_dev_guide/异步任务总述.md→docs/02_dev_guide/03-01_异步任务总述.md+0-0

文件重命名但无更改。

Adocs/02_dev_guide/03-02_Stream管理.md+270-0

@@ -0,0 +1,270 @@
	1	+# Stream管理
	2	+
	3	+## Stream概念
	4	+
	5	+Stream描述了一个在Host下发并在Device上执行的任务队列。
	6	+
	7	+在同一个Stream中，任务按照进入队列的顺序依次执行。当硬件资源充足时，不同Stream上的任务会被调度到不同的硬件资源上并行执行。当硬件资源不足时，不同Stream上的任务可能串行执行。
	8	+
	9	+Stream可以配置优先级、遇错即停、persistent等多种属性。优先级会影响不同Stream上的任务的执行顺序。通常情况下，高优先级Stream上的任务会优先于低优先级Stream上的任务执行。如果Stream配置了persistent属性，则下发在其上的任务不会被立即执行，执行完成后也不会被立即销毁。Persistent Stream适用于模型运行实例构建场景。
	10	+
	11	+相对于主机线程，Stream上的任务是异步执行的。主机线程可以使用Device同步接口来等待当前Context下所有Stream上的任务全部执行完成，或者使用Stream同步接口来等待Stream上的任务全部执行完成。
	12	+
	13	+Runtime中的Stream均为非阻塞式Stream，默认Stream不会跟显式创建的Stream进行隐式同步。
	14	+
	15	+<br>
	16	+<br>
	17	+
	18	+## Stream创建与销毁
	19	+
	20	+调用aclrtCreateStream创建Stream，得到的aclrtStream对象作为后续的内存异步复制、Stream同步、Kernel执行等接口的Stream入参。显式创建的Stream需要调用aclrtDestroyStream接口显式销毁。销毁Stream时，如果Stream上有未完成的任务，则会等待任务完成后再销毁Stream。
	21	+
	22	+以下是创建Stream并在Stream上下发计算任务的代码示例，不可以直接拷贝编译运行，仅供参考。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/blob/master/example/1_basic_features/stream/0_simple_stream)。
	23	+
	24	+```
	25	+// 显式创建一个Stream
	26	+aclrtStream stream;
	27	+aclrtCreateStream(&stream);
	28	+
	29	+// 在Stream上下发Host->Device复制任务、MyKernel任务、和Device->Host复制任务
	30	+aclrtMemcpyAsync(devPtr, devSize, hostPtr, hostSize, ACL_MEMCPY_HOST_TO_DEVICE, stream);
	31	+myKernel<<<8, nullptr, stream>>>();
	32	+aclrtMemcpyAsync(hostPtr, hostSize, devPtr, devSize, ACL_MEMCPY_DEVICE_TO_HOST, stream);
	33	+
	34	+// 销毁Stream（等待Device->Host复制任务执行完成后销毁）
	35	+aclrtDestroyStream(stream);
	36	+```
	37	+
	38	+<br>
	39	+<br>
	40	+
	41	+## 默认Stream
	42	+
	43	+在调用aclrtSetDevice接口或aclrtCreateContext接口时，Runtime会自动创建一个默认Stream。每个Context拥有一个默认Stream。如果不同的Host线程使用相同的Context，它们将共享同一个默认Stream。
	44	+
	45	+对于需要传入Stream参数的API（如aclrtMemcpyAsync），如果使用默认Stream作为入参，则直接传入nullptr。对于没有Stream入参的API（如aclrtMemcpy），则不使用默认Stream。
	46	+
	47	+默认Stream不能显式调用aclrtDestroyStream接口销毁。在调用aclrtResetDevice或aclrtResetDeviceForce接口释放资源时，默认Stream会被自动销毁。
	48	+
	49	+以下是在默认Stream上下发计算任务的代码示例，不可以直接拷贝编译运行，仅供参考。
	50	+
	51	+```
	52	+// 指定Device（接口内部自动创建默认Stream）
	53	+aclrtSetDevice(0);
	54	+
	55	+// 在默认stream上下发Host->Device复制任务、MyKernel任务、和Device->Host复制任务
	56	+aclrtMemcpyAsync(devPtrIn, size, hostPtr, hostSize, ACL_MEMCPY_HOST_TO_DEVICE, nullptr);
	57	+myKernel<<<8, nullptr, nullptr>>>(devPtrIn, devPtrOut, size);
	58	+aclrtMemcpyAsync(hostPtr, hostSize, devPtrOut, size, ACL_MEMCPY_DEVICE_TO_HOST, nullptr);
	59	+
	60	+// 同步默认stream
	61	+aclrtStreamSynchronize(nullptr);
	62	+
	63	+// 复位Device（接口内部自动销毁默认Stream）
	64	+aclrtResetDevice(0);
	65	+```
	66	+
	67	+<br>
	68	+<br>
	69	+
	70	+## 显式同步
	71	+
	72	+对于异步任务接口，主机线程调用异步任务接口后仅代表下发任务，不代表任务执行完成。用户需要显式调用设备同步、流同步等显式同步接口等待任务完成。调用此类显式同步接口后，主机线程会阻塞直到相关的任务执行完成。
	73	+
	74	+### 设备同步：aclrtSynchronizeDevice
	75	+
	76	+阻塞当前主机线程直到当前Device的当前Context中所有显式或隐式创建的Stream完成已下发的所有任务。
	77	+
	78	+以下是设备同步代码示例，不可以直接拷贝编译运行，仅供参考。
	79	+
	80	+```
	81	+// 指定Device
	82	+aclrtSetDevice(0);
	83	+
	84	+// 创建Stream
	85	+aclrtStream stream;
	86	+aclrtCreateStream(&stream);
	87	+
	88	+// 在Stream上下发任务
	89	+......
	90	+
	91	+// 阻塞应用程序运行，直到正在运算中的Device完成运算
	92	+aclrtSynchronizeDevice();
	93	+
	94	+// 资源销毁
	95	+aclrtDestroyStream(stream);
	96	+aclrtResetDevice(0);
	97	+```
	98	+
	99	+### 流同步：aclrtSynchronizeStream
	100	+
	101	+阻塞当前主机线程直到指定的Stream完成已下发的所有任务。
	102	+
	103	+以下是流同步的代码示例，不可以直接拷贝编译运行，仅供参考。
	104	+
	105	+```
	106	+// 创建Stream
	107	+aclrtStream stream;
	108	+aclrtCreateStream(&stream);
	109	+
	110	+// 在Stream上下发任务
	111	+......
	112	+
	113	+// 调用aclrtSynchronizeStream接口，阻塞应用程序运行，直到指定Stream中的所有任务都完成。
	114	+aclrtSynchronizeStream(stream);
	115	+
	116	+// Stream使用结束后，显式销毁Stream
	117	+aclrtDestroyStream(stream);
	118	+```
	119	+
	120	+此外，用户可以使用aclrtStreamQuery查询stream上的任务是否全部执行完成。
	121	+
	122	+
	123	+<br>
	124	+<br>
	125	+
	126	+## Host回调任务
	127	+
	128	+CANN为CPU和NPU之间的异步协作提供了灵活的方式。用户可以使用aclrtLaunchHostFunc在Stream的任意位置插入一个Host回调任务。当本Stream上所有前序任务执行完成后，该Host回调任务会被自动执行，并且会阻塞本Stream上的后续任务执行。
	129	+
	130	+回调函数不能直接或者间接调用CANN Runtime API，否则可能会导致错误或死锁。
	131	+
	132	+以下是在Stream上插入一个Host回调任务的代码示例，不可以直接拷贝编译运行，仅供参考。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/tree/master/example/2_advanced_features/callback/1_callback_hostfunc)。
	133	+
	134	+```
	135	+// Host回调任务
	136	+void myHostCallback(void *args)
	137	+{
	138	+ printf("In MyHostCallback.\n");
	139	+
	140	+ // myKernel1完成后的处理，阻塞MyKernel2的执行
	141	+ ......
	142	+}
	143	+......
	144	+
	145	+// 创建Stream
	146	+aclrtStream stream;
	147	+aclrtCreateStream(&stream);
	148	+
	149	+// 在Stream上下发任务
	150	+aclrtMemcpyAsync(devPtrIn, size, hostPtr, hostSize, ACL_MEMCPY_HOST_TO_DEVICE, stream);
	151	+myKernel1<<<8, nullptr, stream>>>(devPtrIn, devPtrOut, size);
	152	+aclrtLaunchHostFunc(stream, myHostCallback, nullptr);
	153	+myKernel2<<<8, nullptr, stream>>>(devPtrOut, size);
	154	+aclrtMemcpyAsync(hostPtr, hostSize, devPtrOut, size, ACL_MEMCPY_DEVICE_TO_HOST, stream);
	155	+
	156	+// 阻塞应用程序运行，直到指定Stream中的所有任务都完成
	157	+aclrtSynchronizeStream(stream);
	158	+
	159	+// 销毁Stream
	160	+aclrtDestroyStream(stream);
	161	+```
	162	+
	163	+<br>
	164	+<br>
	165	+
	166	+## 配置Stream优先级
	167	+
	168	+在运行时，Device上的调度器会依据各个Stream的优先级来决定任务的执行顺序。高优先级Stream中待执行的任务将优先于低优先级Stream中的任务得到调度，但不会抢占已处于运行状态的低优先级任务。Device在执行过程中不会动态重新评估任务队列，因此提升Stream的优先级不会中断正在执行的任务。
	169	+
	170	+Stream的优先级主要用于影响任务的调度顺序，而非强制规定严格的执行序列。用户可以通过调整Stream的优先级来引导任务的执行顺序，但无法以此强制保证任务间的绝对顺序。
	171	+
	172	+调用aclrtCreateStreamWithConfig接口创建Stream时可指定Stream的优先级。允许设置的优先级范围，可以通过aclrtDeviceGetStreamPriorityRange接口获取最小优先级、最大优先级。
	173	+
	174	+以下为示例代码，不可以直接拷贝编译运行，仅供参考。
	175	+
	176	+```
	177	+// 查询当前设备支持的Stream最小、最大优先级
	178	+aclrtDeviceGetStreamPriorityRange(&leastPriority, &greatestPriority);
	179	+
	180	+// 创建具有最高和最低优先级的Stream
	181	+aclrtStream stream_high;
	182	+aclrtStream stream_low;
	183	+aclrtCreateStreamWithConfig(&stream_high, greatestPriority, ACL_STREAM_FAST_LAUNCH);
	184	+aclrtCreateStreamWithConfig(&stream_low, leastPriority, ACL_STREAM_FAST_LAUNCH);
	185	+```
	186	+
	187	+Stream的优先级在Device范围内生效，而不是在Context范围内生效。
	188	+
	189	+
	190	+<br>
	191	+<br>
	192	+
	193	+## 配置任务遇错即停
	194	+
	195	+CANN支持遇错即停模式（ACL\_STOP\_ON\_FAILURE）和遇错继续模式（ACL\_CONTINUE\_ON\_FAILURE），以支持不同应用对任务执行失败的差异化控制。默认为遇错继续模式。
	196	+
	197	+当Stream上的任务执行失败时，如果配置了遇错即停模式（ACL\_STOP\_ON\_FAILURE），则会停止执行该Context中所有Stream上的任务；如果配置了遇错继续模式（ACL\_CONTINUE\_ON\_FAILURE），则会继续执行Stream上的后续任务。
	198	+
	199	+调用aclrtSetStreamFailureMode接口指定调度模式的示例代码如下，不可以直接拷贝编译运行，仅供参考：
	200	+
	201	+```
	202	+aclrtStream stream;
	203	+aclrtCreateStream(&stream);
	204	+
	205	+// 设置遇错即停模式
	206	+aclrtSetStreamFailureMode(stream, ACL_STOP_ON_FAILURE);
	207	+......
	208	+```
	209	+
	210	+也可以调用aclrtSetStreamAttribute接口指定调度模式的示例代码如下，不可以直接拷贝编译运行，仅供参考：
	211	+
	212	+```
	213	+aclrtStream stream;
	214	+aclrtCreateStream(&stream);
	215	+
	216	+// 设置遇错继续模式
	217	+aclrtSetStreamAttribute(stream, ACL_STREAM_ATTR_FAILURE_MODE, ACL_CONTINUE_ON_FAILURE);
	218	+......
	219	+```
	220	+
	221	+<br>
	222	+<br>
	223	+
	224	+## Persistent流
	225	+
	226	+非Persistent流上的任务在执行完成之后从Stream出队。如果要多次执行某个任务，需要在非Persistent流上多次下发该任务。
	227	+
	228	+Runtime提供了Persistent流支持任务的持久化。在Persistent流上下发的任务不会被立即执行，任务执行完成后也不会被立即销毁。只有在销毁Persistent流时，相关的任务才会被销毁。
	229	+
	230	+调用aclrtCreateStreamWithConfig接口创建Persistent流，Persistent流需要与模型运行实例创建绑定，支持模型的反复执行。以下为示例代码，不可以直接拷贝编译运行，仅供参考：
	231	+
	232	+```
	233	+// 创建Persistent stream
	234	+aclrtStream stream;
	235	+aclrtCreateStreamWithConfig(&stream, 0, ACL_STREAM_PERSISTENT);
	236	+
	237	+// 构建一个模型运行实例
	238	+aclmdlRI modelRI;
	239	+aclmdlRIBuildBegin(&modelRI, 0);
	240	+
	241	+// 把Persistent stream绑定到模型运行实例
	242	+aclmdlRIBindStream(modelRI, stream, ACL_MODEL_STREAM_FLAG_HEAD);
	243	+
	244	+// 在Persistent流上下发任务
	245	+......
	246	+
	247	+// 标记下发任务结束
	248	+aclmdlRIEndTask(modelRI, stream);
	249	+
	250	+// 结束模型运行实例构建
	251	+aclmdlRIBuildEnd(modelRI, nullptr);
	252	+
	253	+// 在默认stream多次执行模型运行实例
	254	+aclmdlRIExecute(modelRI, -1);
	255	+aclmdlRIExecute(modelRI, -1);
	256	+
	257	+// 解除绑定
	258	+aclmdlRIUnbindStream(modelRI, stream);
	259	+
	260	+// 销毁资源
	261	+aclrtDestroyStream(stream);
	262	+aclmdlRIDestroy(modelRI);
	263	+......
	264	+```
	265	+
	266	+
	267	+
	268	+
	269	+
	270	+

Rdocs/02_dev_guide/Kernel加载与执行.md→docs/02_dev_guide/03-03_Kernel加载与执行.md+0-0

文件重命名但无更改。

Rdocs/02_dev_guide/系统任务.md→docs/02_dev_guide/03-04_系统任务.md+0-0

文件重命名但无更改。

Adocs/02_dev_guide/03-05_Event管理.md+186-0

@@ -0,0 +1,186 @@
	1	+# Event管理
	2	+
	3	+## Event概念
	4	+
	5	+Event用于同一Device内、不同Stream之间的任务同步事件。它支持一个任务等待一个事件，例如stream2的任务依赖stream1的任务，想保证stream1中的任务先完成，这时可创建一个Event，将该Event插入到stream1中（Event Record任务），在stream2中插入一个等待Event完成的任务（Event Wait任务）；也支持多个任务等待同一个事件（多等一），例如stream2和stream3中的任务都等待Stream1中的Event完成；同时，Event支持记录事件时间戳信息。
	6	+
	7	+一个任务等待一个事件的图示如下：
	8	+
	9	+![](figures/Event_一个任务等待一个事件.png)
	10	+
	11	+多个任务等待同一个事件的图示如下：
	12	+
	13	+![](figures/Event_多个任务等待一个事件.png)
	14	+
	15	+<br>
	16	+<br>
	17	+
	18	+## Event的创建与销毁
	19	+
	20	+以下是创建两个Event并销毁的代码示例，该示例仅用于说明Event使用方法，不可以直接拷贝编译运行。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/blob/master/example/1_basic_features/event/1_event_timestamp)。
	21	+
	22	+```
	23	+aclrtEvent startEvent;
	24	+aclrtEvent endEvent;
	25	+// 创建Event，接口传入ACL_EVENT_SYNC参数，表示创建的Event用于同步
	26	+aclrtCreateEventExWithFlag(&startEvent, ACL_EVENT_SYNC);
	27	+aclrtCreateEventExWithFlag(&endEvent, ACL_EVENT_SYNC);
	28	+......
	29	+// 销毁Event
	30	+aclrtDestroyEvent(startEvent);
	31	+aclrtDestroyEvent(endEvent);
	32	+```
	33	+
	34	+<br>
	35	+<br>
	36	+
	37	+## Event等待
	38	+
	39	+多Stream之间任务的同步等待可以利用Event实现，例如，若stream2的任务依赖stream1的任务，想保证stream1中的任务先完成，这时可创建一个Event，调用aclrtRecordEvent接口将Event插入到stream1中（通常称为Event Record任务），调用aclrtStreamWaitEvent接口在stream2中插入一个等待Event完成的任务（通常称为Event Wait任务）。
	40	+
	41	+以下为调用aclrtStreamWaitEvent接口的示例代码，不可以直接拷贝编译运行，仅供参考：
	42	+
	43	+```
	44	+// 创建一个Event
	45	+aclrtEvent event;
	46	+aclrtCreateEventExWithFlag(&event, ACL_EVENT_SYNC);
	47	+
	48	+// 创建stream1
	49	+aclrtStream stream1;
	50	+aclrtCreateStream(&stream1);
	51	+
	52	+// 创建stream2
	53	+aclrtStream stream2;
	54	+aclrtCreateStream(&stream2);
	55	+
	56	+// 在stream1上下发任务
	57	+......
	58	+
	59	+// 在stream1末尾添加了一个event
	60	+aclrtRecordEvent(event, stream1);
	61	+
	62	+// 在stream2上下发不依赖stream1执行完成的任务
	63	+......
	64	+
	65	+// 阻塞stream2运行，直到指定event发生，也就是stream1执行完成
	66	+aclrtStreamWaitEvent(stream2, event);
	67	+
	68	+// 在stream2上下发依赖stream1执行完成的任务
	69	+......
	70	+
	71	+// 阻塞应用程序运行，直到stream1和stream2中的所有任务都执行完成
	72	+aclrtSynchronizeStream(stream1);
	73	+aclrtSynchronizeStream(stream2);
	74	+
	75	+// 显式销毁资源
	76	+aclrtDestroyStream(stream1);
	77	+aclrtDestroyStream(stream2);
	78	+aclrtDestroyEvent(event);
	79	+......
	80	+```
	81	+
	82	+<br>
	83	+<br>
	84	+
	85	+## 记录Event时间戳
	86	+
	87	+在[Event的创建与销毁](#Event的创建与销毁)章节中创建的Event可用于统计Stream上计算任务的耗时，代码示例如下。该示例仅用于说明Event使用方法，不可以直接拷贝编译运行。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/blob/master/example/1_basic_features/event/1_event_timestamp)。
	88	+
	89	+```
	90	+uint64_t time = 0;
	91	+float useTime = 0;
	92	+
	93	+// 创建Stream
	94	+aclrtStream stream;
	95	+aclrtCreateStream(&stream);
	96	+
	97	+aclrtEvent startEvent;
	98	+aclrtEvent endEvent;
	99	+// 创建Event，接口传入ACL_EVENT_TIME_LINE参数，表示创建的Event用于记录
	100	+aclrtCreateEventExWithFlag(&startEvent, ACL_EVENT_TIME_LINE);
	101	+aclrtCreateEventExWithFlag(&endEvent, ACL_EVENT_TIME_LINE);
	102	+
	103	+// 插入startEvent
	104	+aclrtRecordEvent(startEvent, stream);
	105	+// 在Stream中下发计算任务
	106	+kernel<<< grid, block, 0, stream>>>(...);
	107	+// 插入endEvent
	108	+aclrtRecordEvent(endEvent, stream);
	109	+aclrtSynchronizeStream(stream);
	110	+
	111	+// 获取时间戳并计算耗时
	112	+aclrtEventElapsedTime(&useTime, startEvent, endEvent);
	113	+```
	114	+
	115	+<br>
	116	+<br>
	117	+
	118	+## Event查询
	119	+
	120	+调用aclrtQueryEventStatus接口可查询指定Event是否完成，非阻塞接口，Event状态包括ACL\_EVENT\_RECORDED\_STATUS\_COMPLETE（所有任务都已经执行完成）、ACL\_EVENT\_RECORDED\_STATUS\_NOT\_READY（有未执行完的任务）。
	121	+
	122	+以下是通过Event实现多线程内存池复用管理机制的代码示例，不可以直接拷贝编译运行，仅供参考。
	123	+
	124	+1. 在A线程中创建内存池，算子所用的内存来源于内存池，在算子后面插入Event Record任务。
	125	+
	126	+ ```
	127	+ // 申请内存池
	128	+ ......
	129	+ // 创建Stream
	130	+ aclrtStream stream;
	131	+ aclrtCreateStream(&stream);
	132	+
	133	+ // 创建Event
	134	+ aclrtEvent event;
	135	+ aclrtCreateEventExWithFlag(&event, ACL_EVENT_CAPTURE_STREAM_PROGRESS);
	136	+
	137	+ // 在Stream中下发计算任务
	138	+ ......
	139	+ // 在算子所在stream插入event
	140	+ aclrtRecordEvent(event, stream);
	141	+ ```
	142	+
	143	+2. 在B线程中调用查询接口，如果查询的Event已经完成，则代表Event Record前面的算子内存都可以被安全的复用。
	144	+
	145	+ ```
	146	+ aclrtEventRecordedStatus status;
	147	+ // 查询线程A的event是否完成
	148	+ aclrtQueryEventStatus(event, &status);
	149	+ if (status == ACL_EVENT_RECORDED_STATUS_COMPLETE) {
	150	+ // event已完成，算子占用的内存可以复用
	151	+ } else {
	152	+ // 算子并未执行完成，该算子占用的内存不能被复用
	153	+ }
	154	+
	155	+ ```
	156	+
	157	+<br>
	158	+<br>
	159	+
	160	+## Event同步
	161	+
	162	+调用aclrtSynchronizeEvent接口阻塞当前主机线程直到指定的Event事件完成。以下为示例代码，不可以直接拷贝编译运行，仅供参考：
	163	+
	164	+```
	165	+// 创建Event
	166	+aclrtEvent event;
	167	+aclrtCreateEventExWithFlag(&event, ACL_EVENT_CAPTURE_STREAM_PROGRESS);
	168	+
	169	+// 创建Stream
	170	+aclrtStream stream;
	171	+aclrtCreateStream(&stream);
	172	+
	173	+// 在Stream上下发任务
	174	+......
	175	+
	176	+// 在Stream上记录Event
	177	+aclrtRecordEvent(event, stream);
	178	+
	179	+// 阻塞应用程序运行直到Event发生
	180	+aclrtSynchronizeEvent(event);
	181	+
	182	+// 显式销毁资源
	183	+aclrtDestroyStream(stream);
	184	+aclrtDestroyEvent(event);
	185	+```
	186	+

Rdocs/02_dev_guide/Notify管理.md→docs/02_dev_guide/03-06_Notify管理.md+0-0

文件重命名但无更改。

Rdocs/02_dev_guide/内存语义同步.md→docs/02_dev_guide/03-07_内存语义同步.md+1-1

	@@ -60,5 +60,5 @@
60	aclrtValueWrite(syncMem, 2, 0, stream2);	60	aclrtValueWrite(syncMem, 2, 0, stream2);
61	```	61	```
62		62
63	-说明：因为内存语义同步机制是基于通用Device内存实现，所以可以通过aclrtMemset/aclrtMemsetAsync初始化和清除同步所用的内存。	63	+说明：因为内存语义同步机制是基于通用Device内存实现，所以可以通过aclrtMemset/aclrtMemsetAsync初始化和清除同步所用的内存。
64		64

Adocs/02_dev_guide/03_异步任务执行.md+10-0

@@ -0,0 +1,10 @@
	1	+# 异步任务执行
	2	+
	3	+- [异步任务总述](03-01_异步任务总述.md)
	4	+- [Stream管理](03-02_Stream管理.md)
	5	+- [Kernel加载与执行](03-03_Kernel加载与执行.md)
	6	+- [系统任务](03-04_系统任务.md)
	7	+- [Event管理](03-05_Event管理.md)
	8	+- [Notify管理](03-06_Notify管理.md)
	9	+- [内存语义同步](03-07_内存语义同步.md)
	10	+

Rdocs/02_dev_guide/单流捕获.md→docs/02_dev_guide/04-01_单流捕获.md+0-0

文件重命名但无更改。

Rdocs/02_dev_guide/跨流捕获.md→docs/02_dev_guide/04-02_跨流捕获.md+0-0

文件重命名但无更改。

Rdocs/02_dev_guide/任务更新.md→docs/02_dev_guide/04-03_任务更新.md+0-0

文件重命名但无更改。

Rdocs/02_dev_guide/ACL-Graph.md→docs/02_dev_guide/04_ACL-Graph.md+0-0

文件重命名但无更改。

Rdocs/02_dev_guide/跨Device的数据交互.md→docs/02_dev_guide/05_多设备编程.md+100-2

	@@ -1,6 +1,103 @@
1	-# 跨Device的数据交互	1	+# 多设备编程
2		2
3	-本节中的“跨Device的数据交互”是指一个进程内、根据硬件组网（例如处于PCIe或者HCCS互联的组网拓扑下）、Device之间能够访问彼此的内存。可以使用aclrtDeviceCanAccessPeer接口查询两个Device之间是否支持数据交互，若支持，再根据访问方向，分别调用aclrtDeviceEnablePeerAccess接口开启一个Device到另一个Device的数据交互功能，例如，调用一次aclrtDeviceEnablePeerAccess接口开启Device 0到Device 1的数据交互，再调用一次aclrtDeviceEnablePeerAccess接口开启Device 1到Device 0的数据交互。若需关闭Device之间的数据交互，可调用aclrtDeviceDisablePeerAccess接口。对于两个进程之间的通信请参见[进程间通信](进程间通信.md)。	3	+## 多设备编程总述
		4	+
		5	+多设备编程使用户能够利用多个NPU（Neural-Network Processing Unit）的综合性能和内存等资源，实现超越单个NPU的性能水平。通常，每个物理NPU在Runtime编程界面中被抽象为一个Device，任务下发时需要有Device中的Context进行支撑，因此多设备编程需要管理多个Device及其相应的Context。
		6	+
		7	+一些常见的多NPU编程方法包括：
		8	+
		9	+- 单个主机线程驱动多个NPU。
		10	+- 多个主机线程，每个线程驱动自己的NPU。
		11	+- 多个单线程主机进程，每个进程驱动自己的NPU。
		12	+- 多个主机进程，每个进程包含多个线程，每个线程驱动自己的NPU。
		13	+
		14	+<br>
		15	+<br>
		16	+
		17	+## 多Device选择
		18	+
		19	+一个Host搭配多Device的场景下，用户可以在Host侧应用程序中通过aclrtGetDeviceCount接口来获取当前Host上搭配的Device数量，Device按照0、1、2、... 的顺序排布。
		20	+
		21	+以下为获取Device信息的代码示例，不可以直接拷贝编译运行，仅供参考：
		22	+
		23	+```
		24	+// 获取Device数量及其对应的属性信息
		25	+uint32_t deviceCount;
		26	+aclrtGetDeviceCount(&deviceCount);
		27	+uint32_t deviceId;
		28	+for (deviceId = 0; deviceId < deviceCount; ++deviceId) {
		29	+ // 按需查询设备属性信息
		30	+ aclrtDevAttr attr = ACL_DEV_ATTR_VECTOR_CORE_NUM;
		31	+ int64_t value;
		32	+ aclrtGetDeviceInfo(deviceId, attr, &value);
		33	+}
		34	+```
		35	+
		36	+此时，可以随时通过aclrtSetDevice接口按线程粒度切换Device（不会影响其他线程）。指定Device后，后续的内存分配、Kernel执行等操作均在该Device上进行，且Stream、Event等也与当前指定的Device相关联。用户可以通过aclrtResetDevice接口释放资源，但更推荐使用aclrtResetDeviceForce接口一次性清理Device上的资源，包括默认Context、默认Stream以及在默认Context下创建的所有Stream。如果默认Context或默认Stream下的任务尚未完成，系统会等待任务完成后才释放资源。在用户程序中，若使用aclrtResetDevice接口，则需确保aclrtSetDevice和aclrtResetDevice接口的调用次数成对出现。
		37	+
		38	+多Device选择的接口调用流程如下图所示：
		39	+
		40	+<img src="figures/同步等待流程_多Device场景.png" width="60%">
		41	+
		42	+以下是多Device选择的代码示例，不可以直接拷贝编译运行，仅供参考：
		43	+
		44	+```
		45	+// 指定Device 0作为计算设备，并将Device 0的默认Context作为当前线程的默认Context
		46	+aclrtSetDevice(0);
		47	+aclrtStream s0;
		48	+aclrtCreateStream(&s0);
		49	+// 执行任务1
		50	+......
		51	+
		52	+// 指定Device 1作为计算设备，并将Device1的默认Context作为当前线程的默认Context
		53	+aclrtSetDevice(1);
		54	+aclrtStream s1;
		55	+aclrtCreateStream(&s1);
		56	+// 执行任务2
		57	+......
		58	+
		59	+// 复位Device 1，释放计算资源，线程默认Context被释放
		60	+// 如需进行继续运行任务，需要显示指定device&context
		61	+aclrtResetDeviceForce(1);
		62	+
		63	+// 复位device0，释放计算资源
		64	+aclrtResetDeviceForce(0);
		65	+```
		66	+
		67	+<br>
		68	+<br>
		69	+
		70	+## Stream和Event的行为
		71	+
		72	+在与当前Device无所属关系的Stream上下发算子将会失败，示例代码如下：
		73	+
		74	+```
		75	+aclrtSetDevice(0); // 指定Device 0作为计算设备
		76	+aclrtStream s0;
		77	+aclrtCreateStream(&s0); // 在Device 0上创建Stream s0
		78	+myKernel<<<8, nullptr, s0>>>(); // 在Device 0上通过Stream s0下发算子
		79	+
		80	+aclrtSetDevice(1); // 指定Device 1作为计算设备
		81	+aclrtStream s1;
		82	+aclrtCreateStream(&s1); // 在Device 1上创建Stream s1
		83	+myKernel<<<8, nullptr, s1>>>(); // 在Device 1上通过Stream s1下发算子
		84	+
		85	+// 算子下发失败
		86	+myKernel<<<8, nullptr, s0>>>(); // 在Device 1上通过Stream s0下发算子
		87	+```
		88	+
		89	+- 当Stream所属的Device和当前操作的Device不相同时，在此Stream上调用aclrtMemcpyAsync会失败。
		90	+- 当Event和Stream关联到不同的Device上时，调用aclrtRecordEvent会失败。
		91	+- 当Event和Stream关联到不同的Device上时，调用aclrtStreamWaitEvent会失败。
		92	+- 当Event所属的Device和当前操作的Device不相同时，aclrtSynchronizeEvent和aclrtQueryEvent会成功。
		93	+
		94	+
		95	+<br>
		96	+<br>
		97	+
		98	+## 跨Device的数据交互
		99	+
		100	+本节中的“跨Device的数据交互”是指一个进程内、根据硬件组网（例如处于PCIe或者HCCS互联的组网拓扑下）、Device之间能够访问彼此的内存。可以使用aclrtDeviceCanAccessPeer接口查询两个Device之间是否支持数据交互，若支持，再根据访问方向，分别调用aclrtDeviceEnablePeerAccess接口开启一个Device到另一个Device的数据交互功能，例如，调用一次aclrtDeviceEnablePeerAccess接口开启Device 0到Device 1的数据交互，再调用一次aclrtDeviceEnablePeerAccess接口开启Device 1到Device 0的数据交互。若需关闭Device之间的数据交互，可调用aclrtDeviceDisablePeerAccess接口。对于两个进程之间的通信请参见[进程间通信](06_进程间通信.md)。
4		101
5	以下是跨Device内存复制的代码示例，不可以直接拷贝编译运行，仅供参考。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/tree/master/example/1_basic_features/device/2_device_P2P)。	102	以下是跨Device内存复制的代码示例，不可以直接拷贝编译运行，仅供参考。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/tree/master/example/1_basic_features/device/2_device_P2P)。
6		103
	@@ -34,3 +131,4 @@ aclFinalize(); // 去初始化
34	```	131	```
35		132
36		133
		134	+

Rdocs/02_dev_guide/进程间通信.md→docs/02_dev_guide/06_进程间通信.md+1-1

	@@ -183,7 +183,7 @@
183		183
184	此处以A、B进程为例，说明一个Device上、两个进程间的物理内存共享的示例代码，不可以直接拷贝编译运行，仅供参考，完整样例代码请参见[Link](https://gitcode.com/cann/runtime/tree/master/example/1_basic_features/memory/8_physical_memory_sharing_withoutpid)。	184	此处以A、B进程为例，说明一个Device上、两个进程间的物理内存共享的示例代码，不可以直接拷贝编译运行，仅供参考，完整样例代码请参见[Link](https://gitcode.com/cann/runtime/tree/master/example/1_basic_features/memory/8_physical_memory_sharing_withoutpid)。
185		185
186	-若A、B进程使用不同的Device，还需配合使用aclrtDeviceEnablePeerAccess接口开启跨Device的数据交互，详细描述请参见[跨Device的数据交互](跨Device的数据交互.md)。	186	+若A、B进程使用不同的Device，还需配合使用aclrtDeviceEnablePeerAccess接口开启跨Device的数据交互，详细描述请参见[跨Device的数据交互](05_多设备编程.md#跨Device的数据交互)。
187		187
188	1. 在A进程中：	188	1. 在A进程中：
189		189

Rdocs/02_dev_guide/运行时核资源控制.md→docs/02_dev_guide/07_运行时核资源控制.md+0-0

文件重命名但无更改。

Rdocs/02_dev_guide/配置AI-Core栈空间大小.md→docs/02_dev_guide/08_配置AI-Core栈空间大小.md+0-0

文件重命名但无更改。

Rdocs/02_dev_guide/获取Runtime错误码.md→docs/02_dev_guide/09_异常处理.md+7-3

	@@ -1,8 +1,10 @@
1	-# 获取Runtime错误码	1	+# 异常处理
2		2
3	-所有Runtime接口都会返回一个错误码。然而，对于异步接口（参见[异步任务执行](异步任务执行.md)），由于接口在Device任务完成前就已经返回，因此无法报告Device上可能发生的异步任务错误。它只能返回在Device任务执行前发生在Host上的错误，例如参数校验失败。如果发生异步任务错误，对应的错误码将在后续某个无关的Runtime接口调用时返回。	3	+## 获取Runtime错误码
4		4
5	-因此，若要在某个异步函数调用后立即检查异步错误，唯一的方法是在调用该函数后立即调用aclrtSynchronizeDevice接口（或使用[显式同步](显式同步.md)中描述的任何其他同步机制）进行同步，并检查aclrtSynchronizeDevice接口返回的错误码。	5	+所有Runtime接口都会返回一个错误码。然而，对于异步接口（参见[异步任务执行](03_异步任务执行.md)），由于接口在Device任务完成前就已经返回，因此无法报告Device上可能发生的异步任务错误。它只能返回在Device任务执行前发生在Host上的错误，例如参数校验失败。如果发生异步任务错误，对应的错误码将在后续某个无关的Runtime接口调用时返回。
		6	+
		7	+因此，若要在某个异步函数调用后立即检查异步错误，唯一的方法是在调用该函数后立即调用aclrtSynchronizeDevice接口（或使用[显式同步](03-02_Stream管理.md#显式同步)中描述的任何其他同步机制）进行同步，并检查aclrtSynchronizeDevice接口返回的错误码。
6		8
7	Runtime会为每个Host线程维护一个错误变量，该变量初始化为ACL\_RT\_SUCCESS，并在每次发生错误（无论是参数校验错误还是异步错误）时被错误码覆盖。aclrtPeekAtLastError接口会返回这个变量的值。aclrtGetLastError接口也会返回这个变量，但同时会将其重置为ACL\_RT\_SUCCESS。	9	Runtime会为每个Host线程维护一个错误变量，该变量初始化为ACL\_RT\_SUCCESS，并在每次发生错误（无论是参数校验错误还是异步错误）时被错误码覆盖。aclrtPeekAtLastError接口会返回这个变量的值。aclrtGetLastError接口也会返回这个变量，但同时会将其重置为ACL\_RT\_SUCCESS。
8		10
	@@ -42,3 +44,5 @@ error = aclrtResetDevice(0);
42		44
43	注意：在遇错继续模式下（具体请参见[配置任务遇错即停](配置任务遇错即停.md)），如果一条Stream上的任务执行出现异常，该Stream上的其他未执行任务仍可继续执行，同时也不会阻止向该Stream及同处于同一Context下的其他Stream下发新任务。此时，aclrtPeekAtLastError和aclrtGetLastError返回的可能不是首次错误的信息。	45	注意：在遇错继续模式下（具体请参见[配置任务遇错即停](配置任务遇错即停.md)），如果一条Stream上的任务执行出现异常，该Stream上的其他未执行任务仍可继续执行，同时也不会阻止向该Stream及同处于同一Context下的其他Stream下发新任务。此时，aclrtPeekAtLastError和aclrtGetLastError返回的可能不是首次错误的信息。
44		46
		47	+
		48	+

Ddocs/02_dev_guide/Device内存使用.md+0-73

	@@ -1,73 +0,0 @@
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	-

Ddocs/02_dev_guide/Event同步.md+0-27

	@@ -1,27 +0,0 @@
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	-

Ddocs/02_dev_guide/Event查询.md+0-39

	@@ -1,39 +0,0 @@
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	-

Ddocs/02_dev_guide/Event概念.md+0-12

	@@ -1,12 +0,0 @@
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	-

Ddocs/02_dev_guide/Event的创建与销毁.md+0-17

	@@ -1,17 +0,0 @@
1	-# Event的创建与销毁
2	-
3	-以下是创建两个Event并销毁的代码示例，该示例仅用于说明Event使用方法，不可以直接拷贝编译运行。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/blob/master/example/1_basic_features/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	-
17	-

Ddocs/02_dev_guide/Event等待.md+0-45

	@@ -1,45 +0,0 @@
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	-

Ddocs/02_dev_guide/Event管理.md+0-14

	@@ -1,14 +0,0 @@
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	-

Ddocs/02_dev_guide/Host回调任务.md+0-38

	@@ -1,38 +0,0 @@
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/2_advanced_features/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	-
38	-

Ddocs/02_dev_guide/Host锁页内存使用.md+0-68

	@@ -1,68 +0,0 @@
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	-

Ddocs/02_dev_guide/Persistent流.md+0-42

	@@ -1,42 +0,0 @@
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	-

Ddocs/02_dev_guide/Stream创建与销毁.md+0-21

	@@ -1,21 +0,0 @@
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/1_basic_features/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	-
21	-

Ddocs/02_dev_guide/Stream和Event的行为.md+0-24

	@@ -1,24 +0,0 @@
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	-

Ddocs/02_dev_guide/Stream概念.md+0-12

	@@ -1,12 +0,0 @@
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	-

Ddocs/02_dev_guide/Stream管理.md+0-18

	@@ -1,18 +0,0 @@
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	-

Ddocs/02_dev_guide/dev_guide.md+0-51

	@@ -1,51 +0,0 @@
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	-

Ddocs/02_dev_guide/内存管理.md+0-10

	@@ -1,10 +0,0 @@
1	-# 内存管理
2	-
3	-- [内存管理总述](内存管理总述.md)
4	-
5	-- [Device内存使用](Device内存使用.md)
6	-
7	-- [Host锁页内存使用](Host锁页内存使用.md)
8	-
9	-- [虚拟内存管理](虚拟内存管理.md)
10	-

Ddocs/02_dev_guide/内存管理总述.md+0-9

	@@ -1,9 +0,0 @@
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	-

Ddocs/02_dev_guide/多Device选择.md+0-50

	@@ -1,50 +0,0 @@
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	-

Ddocs/02_dev_guide/多设备编程.md+0-10

	@@ -1,10 +0,0 @@
1	-# 多设备编程
2	-
3	-- [多设备编程总述](多设备编程总述.md)
4	-
5	-- [多Device选择](多Device选择.md)
6	-
7	-- [Stream和Event的行为](Stream和Event的行为.md)
8	-
9	-- [跨Device的数据交互](跨Device的数据交互.md)
10	-

Ddocs/02_dev_guide/多设备编程总述.md+0-11

	@@ -1,11 +0,0 @@
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	-

Ddocs/02_dev_guide/异常处理.md+0-4

	@@ -1,4 +0,0 @@
1	-# 异常处理
2	-
3	-- [获取Runtime错误码](获取Runtime错误码.md)
4	-

Ddocs/02_dev_guide/异步任务执行.md+0-16

	@@ -1,16 +0,0 @@
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	-

Ddocs/02_dev_guide/显式同步.md+0-52

	@@ -1,52 +0,0 @@
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	-

Ddocs/02_dev_guide/记录Event时间戳.md+0-31

	@@ -1,31 +0,0 @@
1	-# 记录Event时间戳
2	-
3	-在[Event的创建与销毁](Event的创建与销毁.md)章节中创建的Event可用于统计Stream上计算任务的耗时，代码示例如下。该示例仅用于说明Event使用方法，不可以直接拷贝编译运行。完整样例代码请参见[Link](https://gitcode.com/cann/runtime/blob/master/example/1_basic_features/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	-
31	-

Ddocs/02_dev_guide/配置Stream优先级.md+0-23

	@@ -1,23 +0,0 @@
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	-

Ddocs/02_dev_guide/配置任务遇错即停.md+0-28

	@@ -1,28 +0,0 @@
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	-

Ddocs/02_dev_guide/默认Stream.md+0-26

	@@ -1,26 +0,0 @@
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	-

Mdocs/README.md+1-1

	@@ -3,6 +3,6 @@ Runtime文档包括快速入门、编程指南、API参考部分的内容，如
3		3
4	⦁ [快速入门](01_quick_start/quick_start.md)：介绍Runtime的基本功能、编程模型以及关键概念等。	4	⦁ [快速入门](01_quick_start/quick_start.md)：介绍Runtime的基本功能、编程模型以及关键概念等。
5		5
6	-⦁ [编程指南](02_dev_guide/dev_guide.md)：介绍了如何基于Runtime API进行初始化、内存管理、异步任务执行、ACL Graph、多设备编程等。	6	+⦁ [编程指南](02_dev_guide/00_dev_guide.md)：介绍了如何基于Runtime API进行初始化、内存管理、异步任务执行、ACL Graph、多设备编程等。
7		7
8	⦁ [API参考](03_api_ref/api_ref.md)：介绍了Runtime API的功能、参数及使用说明等。	8	⦁ [API参考](03_api_ref/api_ref.md)：介绍了Runtime API的功能、参数及使用说明等。

@@ -0,0 +1,10 @@
	1	+# 异步任务执行
	2	+
	3	+- [异步任务总述](03-01_异步任务总述.md)
	4	+- [Stream管理](03-02_Stream管理.md)
	5	+- [Kernel加载与执行](03-03_Kernel加载与执行.md)
	6	+- [系统任务](03-04_系统任务.md)
	7	+- [Event管理](03-05_Event管理.md)
	8	+- [Notify管理](03-06_Notify管理.md)
	9	+- [内存语义同步](03-07_内存语义同步.md)
	10	+

	@@ -1,10 +0,0 @@
1	-# 内存管理
2	-
3	-- [内存管理总述](内存管理总述.md)
4	-
5	-- [Device内存使用](Device内存使用.md)
6	-
7	-- [Host锁页内存使用](Host锁页内存使用.md)
8	-
9	-- [虚拟内存管理](虚拟内存管理.md)
10	-

	@@ -1,10 +0,0 @@
1	-# 多设备编程
2	-
3	-- [多设备编程总述](多设备编程总述.md)
4	-
5	-- [多Device选择](多Device选择.md)
6	-
7	-- [Stream和Event的行为](Stream和Event的行为.md)
8	-
9	-- [跨Device的数据交互](跨Device的数据交互.md)
10	-