昇腾社区首页
开发者
资源
支持
积分兑换 NEW

编程模型

主机、设备编程模式

主机（或称Host）：指X86服务器CPU、ARM服务器CPU，通过总线与一个或多个设备互联，利用设备提供的NN（Neural-Network）等计算能力完成业务。

设备（或称Device、NPU）：指安装了AI处理器的硬件，通过总线（例如PCIe、HCCS等）与主机相连，提供NN等计算能力。HCCS是Huawei Cache Coherence System，表示华为缓存一致性系统。

主机与设备的关系如下图所示：

总结如下：

  1. 主机和设备各自拥有独立的内存空间。Runtime提供了主机侧和设备侧的内存申请接口，以及主机与设备之间的内存复制接口。在编程时，需要区分主机和设备的内存申请，并显式调用内存复制接口，将数据从主机侧复制到设备侧，以使设备硬件加速器在访问本地内存时达到最佳性能。
  2. 主机与设备之间采用异步并行的执行方式。

    主机将任务（或称Task）下发到设备后，不会等待设备任务执行完成就立即返回；设备随即开始调度并执行下发的任务；主机侧的CPU可以与设备侧的加速器并行工作。异步任务下发的接口通常会带有Stream参数，表示将任务下发到对应的Stream中执行。

    当主机需要获取设备的计算结果时，必须发起显式的同步API调用，同步API会阻塞主机CPU，直到设备侧任务执行完成才返回。

    通过这种异步并行执行机制，可以有效隐藏主机处理时间或主机与设备间的数据传输延迟，提高吞吐量，缩短端到端的执行时间。

  3. 异步任务下发到Stream中任务的执行方式

    Stream中的任务保序执行，Stream间的任务并行执行。例如下图中主机侧顺序启动Kernel1、Kernel2和Kernel3任务，具体执行顺序如下：

    • Kernel1和Kernel3位于同一个Stream中，因此Kernel3需要等待Kernel1执行完毕才能开始执行。
    • Kernel2与Kernel1、Kernel3不在同一个Stream中，因此Kernel2可以与Kernel1、Kernel3并行执行。

典型执行流程

基于Runtime编程的典型执行流程图如下所示：

接口调用

处理

int32_t devId=0;

aclrtSetDevice(devId);

Device 0初始化，主要包括以下操作：

1. 创建并初始化Device对象

2. 为Device创建默认Context;

3. 为默认Context创建默认流;

4. 启动Device侧的CPU执行器进程。

aclrtStream stream1;

aclrtCreateStream(&stream1);

在当前Context下创建Stream：

1. 调用驱动创建任务队列;

2. Runtime侧创建Stream对象，Stream与任务队列关联。

3. 将Stream纳入到当前Context中管理。

uint64_t size=1024;

void *hostPtr=nullptr;

aclrtMallocHost(&hostPtr, size);

申请主机内存。

void *devPtr=nullptr;

aclrtMalloc(&devPtr, size);

申请设备内存。

aclrtMemcpy(devPtr, size, hostPtr, size, ACL_MEMCPY_HOST_TO_DEVICE)

Host到Device的内存同步拷贝，主要包括以下操作：

1. 根据源、目的地址构造DMA（Direct Memory Access）描述符后下发DMA任务。

2. 等待DMA任务完成，接口返回。

myKernel<<<numBlocks, nullptr, stream1>>>(devPtr);
  1. 向stream1下发myKernel计算任务，numBlocks用于指定计算任务的block数量，详细操作如下：
    1. myKernel的二进制代码注册并加载至Device （仅首次）。
    2. 将myKernel的执行参数拷贝至Device。
    3. 分配任务管理资源，并创建任务描述符。
    4. 下发任务到此流对应的任务队列。
    5. 下发完任务即接口异步返回。
  2. 任务调度器开始调度，详细操作如下：
    1. 调度器按任务队列的优先级采用绝对优先级（SP）方式调度。
    2. 解析任务队列中下发的任务描述符。
    3. 根据任务类型判断对应加速单元是否有空闲。
    4. 有空闲且满足任务的核数量要求，则将任务调度给对应加速单元（例如，CPU算子会调度给CPU执行器）。
    5. 加速单元执行任务。
    6. 调度器等待加速单元执行完成，刷新任务状态和加速单元忙闲状态。当任务执行异常，调度器会处理异常。
    7. 继续调度其他任务。

aclrtSynchronizeStream(stream1);

Stream同步，详细操作如下：

1. 接口会阻塞当前CPU线程。

2. 通过轮循+中断机制同步流对应的任务队列执行状态。

3. 流上任务全部完成，此接口返回。

aclrtMemcpy(hostPtr, size, devPtr, size, ACL_MEMCPY_DEVICE_TO_HOST);

Device到Host的内存同步拷贝，将结果复制回Host：

1. 根据源、目的地址构造DMA描述符后下发DMA任务。

2. 等待DMA任务完成，接口返回。

aclrtFree(devPtr);

aclrtFreeHost(hostPtr);

释放设备内存。

释放主机内存。

aclrtResetDeviceForce(devId);

释放Device。

Runtime主要编程概念

  • Host, 是Runtime对主机的抽象。
  • 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之间的关联和切换过程，仅供参考，不可以直接拷贝编译运行。
 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
// 初始时，线程未关联任何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(&current_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。

示例代码如下，仅供参考，不可以直接拷贝编译运行：

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
// ......
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生命周期也终止

编写高性能应用程序的建议

遵循如下基本原则：

  1. 主机侧与Device侧执行异步执行。主机侧要能及时下发足够任务至Device，确保加速硬件始终处于计算状态。
  2. 采用多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的创建按照算子执行硬件划分。
父主题： 快速入门