CANN Runtime编程模型详解
Runtime编程模型
【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime
主机、设备编程模式
主机（或称Host）：指X86服务器CPU、ARM服务器CPU，通过总线与一个或多个设备互联，利用设备提供的NN（Neural-Network）等计算能力完成业务。
设备（或称Device、NPU）：指安装了AI处理器的硬件，通过总线（例如PCIe、HCCS等）与主机相连，提供NN等计算能力。HCCS是Huawei Cache Coherence System，表示华为缓存一致性系统。
主机与设备的关系如下图所示：
总结如下：
**主机和设备各自拥有独立的内存空间。**Runtime提供了主机侧和设备侧的内存申请接口，以及主机与设备之间的内存复制接口。在编程时，需要区分主机和设备的内存申请，并显式调用内存复制接口，将数据从主机侧复制到设备侧，以使设备硬件加速器在访问本地内存时达到最佳性能。
主机与设备之间采用异步并行的执行方式。
主机将任务（或称Task）下发到设备后，不会等待设备任务执行完成就立即返回；设备随即开始调度并执行下发的任务；主机侧的CPU可以与设备侧的加速器并行工作。异步任务下发的接口通常会带有Stream参数，表示将任务下发到对应的Stream中执行。
当主机需要获取设备的计算结果时，必须发起显式的同步API调用，同步API会阻塞主机CPU，直到设备侧任务执行完成才返回。
通过这种异步并行执行机制，可以有效隐藏主机处理时间或主机与设备间的数据传输延迟，提高吞吐量，缩短端到端的执行时间。
异步任务下发到Stream中任务的执行方式
Stream中的任务保序执行，Stream间的任务并行执行。例如下图中主机侧顺序启动Kernel1、Kernel2和Kernel3任务，具体执行顺序如下：
- Kernel1和Kernel3位于同一个Stream中，因此Kernel3需要等待Kernel1执行完毕才能开始执行。
- Kernel2与Kernel1、Kernel3不在同一个Stream中，因此Kernel2可以与Kernel1、Kernel3并行执行。
典型执行流程
基于Runtime编程的典型执行流程图如下所示：
Device初始化，以Device 0为例。
int32_t devId=0; aclrtSetDevice(devId);接口内部涉及如下操作：
- 创建并初始化Device对象。
- 为Device创建默认Context。
- 为默认Context创建默认流。
- 启动Device侧的CPU执行器进程。
在当前Context下创建Stream。
aclrtStream stream1; aclrtCreateStream(&stream1);接口内部涉及如下操作：
- 调用驱动创建任务队列。
- Runtime侧创建Stream对象，Stream与任务队列关联。
- 将Stream纳入到当前Context中管理。
申请主机内存。
uint64_t size=1024; void *hostPtr=nullptr; aclrtMallocHost(&hostPtr, size);申请设备内存。
void *devPtr=nullptr; aclrtMalloc(&devPtr, size);Host到Device的内存同步拷贝。
aclrtMemcpy(devPtr, size, hostPtr, size, ACL_MEMCPY_HOST_TO_DEVICE);Host到Device的内存同步拷贝，主要包括以下操作：
- 根据源、目的地址构造DMA（Direct Memory Access）描述符后下发DMA任务。
- 等待DMA任务完成，接口返回。
下发myKernel计算任务。
myKernel<<<numBlocks, nullptr, stream1>>>(devPtr);下发计算任务及任务调度的详细步骤如下：
向stream1下发myKernel计算任务，numBlocks用于指定计算任务的block数量，详细操作如下：
(a) myKernel的二进制代码注册并加载至Device （仅首次）。
(b) 将myKernel的执行参数拷贝至Device。
(c) 分配任务管理资源，并创建任务描述符。
(d) 下发任务到此流对应的任务队列。
(e) 下发完任务即接口异步返回。
任务调度器开始调度，详细操作如下： (a) 调度器按任务队列的优先级采用绝对优先级（SP）方式调度。
(b) 解析任务队列中下发的任务描述符。
(c) 根据任务类型判断对应加速单元是否有空闲。
(d) 有空闲且满足任务的核数量要求，则将任务调度给对应加速单元（例如，CPU算子会调度给CPU执行器）。
(e) 加速单元执行任务。
(f) 调度器等待加速单元执行完成，刷新任务状态和加速单元忙闲状态。当任务执行异常，调度器会处理异常。
(g) 继续调度其他任务。
Stream同步。
aclrtSynchronizeStream(stream1);Stream同步，详细操作如下：
- 接口会阻塞当前CPU线程。
- 通过轮循+中断机制同步流对应的任务队列执行状态。
- 流上任务全部完成，此接口返回。
Device到Host的内存同步拷贝，将结果复制回Host。
aclrtMemcpy(hostPtr, size, devPtr, size, ACL_MEMCPY_DEVICE_TO_HOST);根据源、目的地址构造DMA描述符后下发DMA任务。等待DMA任务完成，接口返回。
释放设备核主机内存。
aclrtFree(devPtr); aclrtFreeHost(hostPtr);释放Device资源。
aclrtResetDeviceForce(devId);
Runtime主要编程概念
- Host, 是Runime对主机的抽象。
- Device，是对AI处理器所属设备的抽象，通常Host与Device关系为1：N。用户APP可以调用acl接口，例如aclrtSetDevice，指定当前用于运算的硬件设备。
- Context，是Device的逻辑运行环境，Context与Device的关系为N：1，即每个Context必定隶属于一个唯一的Device。Context负责管理运行资源对象（包括Stream、Event和Notify，但不包括内存）的生命周期；不同Context中的对象是完全隔离的，例如，不同Context的Stream和Event是完全隔离的，无法建立同步等待关系；运行出错同样按Context隔离。
- Stream，是Device提供的逻辑任务执行队列，可以异步地向Stream中添加任务，在同一个Stream中的任务会严格按FIFO方式执行。Stream与Context的关系是N：1，某条Stream一定属于唯一的Context。
- Task，可被添加到Stream中的执行任务，可以分计算类任务、内存拷贝、事件同步类任务。Task与Stream的关系是N：1，某个Task会被加入到唯一的Stream。
Device、Context、Stream之间的关系如下图所示：
线程关联Context
Runtime的大多数API接口没有device id参数，因为这些API接口所作用的Device是从调用线程关联的Context中获取的。因此，当主机线程调用Runtime API时，要遵循如下要求：
- 线程（主机侧的CPU线程）要关联Context后，才能正确调用Runtime API。
- 线程同一时刻只能关联一个Context。
- 应用程序可以显式创建Context来达成运行资源隔离的业务诉求。此场景中，同一进程内Context可被所有线程可见，线程可以通过aclrtGetCurrentContext，aclrtSetCurrentContext进行切换Context。
以下示例说明了线程在调用Context相关接口时，线程与Context之间的关联和切换过程，仅供参考，不可以直接拷贝编译运行。
// 初始时，线程未关联任何Context aclInit(nullptr); aclrtSetDevice(0); // aclrtSetDevice会创建默认Context，同时将线程关联默认Context // 线程关联的Context: 默认context aclrtContext ctx1, ctx2, current_ctx; aclrtCreateContext(&ctx1, 0); // Device 0显式创建ctx1, 会将线程关联ctx1 // 线程关联的Context: ctx1 aclrtGetCurrentContext(¤t_ctx); // 获取当前线程关联的Context, 此时返回的current_ctx==ctx1 // 线程关联的Context: ctx1 aclrtCreateContext(&ctx2, 0); // Device 0又显式创建ctx2, 会将线程关联ctx2 // 线程关联的Context: ctx2 aclrtSetCurrentContext(current_ctx); // 切换Context，由于current_ctx=ctx1，线程关联ctx1 // 线程关联的context: ctx1 aclrtSetCurrentContext(ctx2); // 切换Context // 线程关联的context: ctx2 ..... aclrtDestroyContext(ctx2); // 当前线程正关联ctx2，销毁ctx2时会同时将线程也ctx2去关联 // 线程关联的context: NA // ctx2已销毁，应该切换到其他ctx(如ctx1) aclrtSetCurrentContext(ctx1); // 线程关联的context: ctx1 ..... aclrtDestroyContext(ctx1); // 当前线程正关联ctx1，销毁ctx1时会同时将线程也ctx1去关联 // 线程关联的context: NA aclrtResetDeviceForce(0);默认Context和默认Stream的使用场景
Device上执行操作下发前，必须有Context和Stream，这个Context、Stream可以显式创建，也可以隐式创建。隐式创建的Context、Stream就是默认Context、默认Stream。
默认Stream作为接口入参时，直接传NULL。
默认Context不允许用户执行aclrtGetCurrentContext或aclrtSetCurrentContext操作，也不允许执行aclrtDestroyContext操作。
默认Context、默认Stream一般适用于简单应用，用户仅需要一个Device的计算场景下。多线程应用程序建议使用显式创建的Context和Stream。
示例代码如下，仅供参考，不可以直接拷贝编译运行：
// ...... uint32_t numBlocks = 32; uint64_t size = 1024; void *devPtr = nullptr; aclInit(nullptr); aclrtSetDevice(0); /* 已经创建了一个默认Context，在默认Context中创建了一个默认Stream，并且在当前线程可用 */ ...... aclrtMalloc(&devPtr, size ); myKernel<<<numBlocks, nullptr, nullptr>>>(devPtr); // <<< >>>中第三参数nullptr表示在默认Stream上执行 aclrtSynchronizeStream(nullptr); /* 等待计算任务全部完成，用户根据需要获取计算任务的输出结果 */ ...... aclrtResetDeviceForce(0); // 释放Device 0，对应的默认Context及默认Stream生命周期也终止编写高性能应用程序的建议
遵循如下基本原则：
主机侧与Device侧执行异步执行。主机侧要能及时下发足够任务至Device，确保加速硬件始终处于计算状态。
采用多Stream方式充分利用Device上不同种类硬件加速器实现并发执行。 如下图所示，CANN Runtime可以协同调度多种硬件加速器，不同代AI处理器支持的硬件加速器不同，需以实际硬件用户手册中的说明为准。
推荐如下方式：
- 单线程中创建并使用多个Stream。如果单线程的性能足以满足向多个Stream下发任务，以充分利用Device的算力，建议采用单线程模式。
- 当单线程性能不足时，可以采用多线程模式来提升主机侧任务下发的性能。推荐每个线程创建并使用各自的Stream下发任务；不推荐多个线程并发向同一个Stream下发任务，这将引入锁操作，且多个线程间下发的任务是乱序的。
- Stream上下发的单个任务占不满AI Core时，可以使用多Stream下发可并行执行的任务来充分利用AI Core资源。
- AI处理器中包含多种硬件加速器，例如AI Core、AI CPU、DVPP（Digital Vision Pre-Processing）、Random（随机数生成器）等，这些硬件加速器对应不同类型的任务，建议多Stream的创建按照算子执行硬件划分。
【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime
创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考