1. MXMACA运行时简介

MXMACA^® 运行时（Runtime）API提供在主机上执行的C和C++函数，用于分配和释放设备内存，在主机内存和设备内存之间传输数据，以及管理具有多个设备的系统。本文档介绍了MXMACA驱动提供的所有运行时API。

1.1. 常规行为-流同步

1.1.1. 默认流

当 0 作为 mcStream_t 传递或由隐式操作流的API传递时，使用的默认流可以配置为具有遗留（legacy）同步行为或每线程（per-thread）同步行为，详细信息请参见 1.1.1.1 遗留默认流和 1.1.1.2 每线程默认流。

可以使用以下任一选项控制行为。

默认流的行为有两种控制方法。

编译时使用mxcc选项 --default-stream：

mxcc –x maca --default-stream {legacy|null|per-thread} –my_code.cpp –o my_app

在包含任意MXMACA头文件之前，可通过定义MACA_API_Per_thread_DEFAULT_STREAM宏来启用每线程行为：
```
#define MACA_API_PER_THREAD_DEFAULT_STREAM
#include <mc_runtime_api.h>
```

1.1.1.1. 遗留默认流

遗留默认流是一个隐式流，它与同一个 mcContext 中除非阻塞流之外的所有流同步，如下所述。对于仅使用运行时API的应用程序，每个设备将有一个上下文。当在遗留流中执行操作（如核函数启动或 mcStreamWaitEvent ）时，遗留流先等待所有阻塞流，该操作在遗留流中排队进行，然后所有阻塞流等待遗留流。

例如，运行以下代码，先在流 my_stream 中启动内核 my_kernel_1，在遗留流中启动内核 my_kernel_2，然后在流 my_stream 中启动内核 my_kernel_3：

my_kernel_1<<<1, 1, 0, my_stream >>>();
my_kernel_2<<<1, 1>>>();
my_kernel_3<<<1, 1, 0, my_stream >>>()

产生的结果行为是：my_kernel_2 会阻塞 my_kernel_1，且 my_kernel_3 会阻塞 my_kernel_2。

可以使用 mcStreamNonBlocking 标志和流创建API来创建不与遗留流同步的非阻塞流。

可以通过 mcStream_t 句柄 mcStreamLegacy 显式使用遗留默认流。

1.1.1.2. 每线程默认流

每线程默认流是线程和 mcContext 本地的隐式流，不与其他流同步（正如显式创建的流）。每线程默认流不是非阻塞流，如果在程序中使用遗留默认流和每线程默认流，它将与遗留默认流同步。

可以通过 mcStream_t 句柄 mcStreamPerThread 显式使用每线程默认流。

1.2. 常规行为-API同步

运行时API提供同步和异步形式的memcpy和memset函数，后者具有Async后缀。因为每个函数都可能会表现出同步或异步行为，具体取决于传递给函数的参数。

1.2.1. Memcpy

在运行时API参考文档中，每个memcpy函数都被分为同步或异步，定义如下。

1.2.1.1. 同步

所有涉及同一内存区域的传输相对于主机都是完全同步的。
对于从分页（pageable）主机内存到设备内存的传输，启动复制前执行流同步。一旦分页缓冲区被复制到分级存储，用于到设备内存的DMA传输，该函数就会返回，但DMA到目标内存的传输可能尚未完成。
对于从固页（pinned）主机内存到设备内存的传输，该函数相对于主机是同步的。
对于从设备内存到分页或固页主机内存的传输，该函数仅在复制完成后返回。
对于从设备内存到设备内存的传输，不执行主机端同步。
对于从任意主机内存到任意主机内存的传输，该函数相对于主机是完全同步的。

1.2.1.2. 异步

对于从设备内存到分页主机内存的传输，该函数仅在复制完成后返回。
对于从任意主机内存到任意主机内存的传输，该函数相对于主机是完全同步的。
如果必须先将分页内存暂存到固页内存，MXMACA驱动可以与流同步，并将副本暂存到固页内存。
对于其他所有传输，该函数是完全异步的。

1.2.2. Memset

同步memset函数相对于主机是异步的，除非目标是固页主机内存或统一内存区域，在这种情况下，它们是完全同步的。异步版本相对于主机始终是异步的。

1.2.3. 核函数启动

核函数启动相对于主机是异步的。有关并发内核执行和数据传输的详细信息，参见《曦云^® 系列通用GPU 运行时API编程指南》。