多设备编程
多设备编程总述
多设备编程使用户能够利用多个NPU（Neural-Network Processing Unit）的综合性能和内存等资源，实现超越单个NPU的性能水平。通常，每个物理NPU在Runtime编程界面中被抽象为一个Device，任务下发时需要有Device中的Context进行支撑，因此多设备编程需要管理多个Device及其相应的Context。
一些常见的多NPU编程方法包括：
- 单个主机线程驱动多个NPU。
- 多个主机线程，每个线程驱动自己的NPU。
- 多个单线程主机进程，每个进程驱动自己的NPU。
- 多个主机进程，每个进程包含多个线程，每个线程驱动自己的NPU。
多Device选择
一个Host搭配多Device的场景下，用户可以在Host侧应用程序中通过aclrtGetDeviceCount接口来获取当前Host上搭配的Device数量，Device按照0、1、2、... 的顺序排布。
以下为获取Device信息的代码示例，不可以直接拷贝编译运行，仅供参考：
// 获取Device数量及其对应的属性信息
uint32_t deviceCount;
aclrtGetDeviceCount(&deviceCount);
uint32_t deviceId;
for (deviceId = 0; deviceId < deviceCount; ++deviceId) {
// 按需查询设备属性信息
aclrtDevAttr attr = ACL_DEV_ATTR_VECTOR_CORE_NUM;
int64_t value;
aclrtGetDeviceInfo(deviceId, attr, &value);
}
此时，可以随时通过aclrtSetDevice接口按线程粒度切换Device（不会影响其他线程）。指定Device后，后续的内存分配、Kernel执行等操作均在该Device上进行，且Stream、Event等也与当前指定的Device相关联。用户可以通过aclrtResetDevice接口释放资源，但更推荐使用aclrtResetDeviceForce接口一次性清理Device上的资源，包括默认Context、默认Stream以及在默认Context下创建的所有Stream。如果默认Context或默认Stream下的任务尚未完成，系统会等待任务完成后才释放资源。在用户程序中，若使用aclrtResetDevice接口，则需确保aclrtSetDevice和aclrtResetDevice接口的调用次数成对出现。
多Device选择的接口调用流程如下图所示：
以下是多Device选择的代码示例，不可以直接拷贝编译运行，仅供参考：
// 指定Device 0作为计算设备，并将Device 0的默认Context作为当前线程的默认Context
aclrtSetDevice(0);
aclrtStream s0;
aclrtCreateStream(&s0);
// 执行任务1
......
// 指定Device 1作为计算设备，并将Device1的默认Context作为当前线程的默认Context
aclrtSetDevice(1);
aclrtStream s1;
aclrtCreateStream(&s1);
// 执行任务2
......
// 复位Device 1，释放计算资源，线程默认Context被释放
// 如需继续运行任务，需要显式指定Device和Context
aclrtResetDeviceForce(1);
// 复位device0，释放计算资源
aclrtResetDeviceForce(0);
Stream和Event的行为
在与当前Device无所属关系的Stream上下发算子将会失败，示例代码如下：
aclrtSetDevice(0); // 指定Device 0作为计算设备
aclrtStream s0;
aclrtCreateStream(&s0); // 在Device 0上创建Stream s0
myKernel<<<8, nullptr, s0>>>(); // 在Device 0上通过Stream s0下发算子
aclrtSetDevice(1); // 指定Device 1作为计算设备
aclrtStream s1;
aclrtCreateStream(&s1); // 在Device 1上创建Stream s1
myKernel<<<8, nullptr, s1>>>(); // 在Device 1上通过Stream s1下发算子
// 算子下发失败
myKernel<<<8, nullptr, s0>>>(); // 在Device 1上通过Stream s0下发算子
- 当Stream所属的Device和当前操作的Device不相同时，在此Stream上调用aclrtMemcpyAsync会失败。
- 当Event和Stream关联到不同的Device上时，调用aclrtRecordEvent会失败。
- 当Event和Stream关联到不同的Device上时，调用aclrtStreamWaitEvent会失败。
- 当Event所属的Device和当前操作的Device不相同时，aclrtSynchronizeEvent和aclrtQueryEvent会成功。
跨Device的数据交互
本节中的“跨Device的数据交互”是指一个进程内、根据硬件组网（例如处于PCIe或者HCCS互联的组网拓扑下）、Device之间能够访问彼此的内存。可以使用aclrtDeviceCanAccessPeer接口查询两个Device之间是否支持数据交互，若支持，再根据访问方向，分别调用aclrtDeviceEnablePeerAccess接口开启一个Device到另一个Device的数据交互功能，例如，调用一次aclrtDeviceEnablePeerAccess接口开启Device 0到Device 1的数据交互，再调用一次aclrtDeviceEnablePeerAccess接口开启Device 1到Device 0的数据交互。若需关闭Device之间的数据交互，可调用aclrtDeviceDisablePeerAccess接口。对于两个进程之间的通信请参见进程间通信。
以下是跨Device内存复制的代码示例，不可以直接拷贝编译运行，仅供参考。完整样例代码请参见Link。
aclInit(NULL); // 初始化
int32_t canAccessPeer = 0;
aclrtDeviceCanAccessPeer(&canAccessPeer, 0, 1); // 查询Device 0和Device 1之间是否支持数据交互
if (canAccessPeer == 1) {
aclrtSetDevice(0); // Device 0下的操作
uint32_t reserveFlag = 0U;
aclrtDeviceEnablePeerAccess(1, reserveFlag); // 开启当前Device(Device 0)到指定Device(Device 1）的数据交互
void *dev0Mem = nullptr;
aclrtMalloc(&dev0Mem, 10, ACL_MEM_MALLOC_HUGE_FIRST_P2P);
aclrtSetDevice(1); // Device 1下的操作
aclrtDeviceEnablePeerAccess(0, reserveFlag); // 开启当前Device(Device 1)到指定Device(Device 0）的数据交互
void *dev1Mem = nullptr;
aclrtMalloc(&dev1Mem, 10, ACL_MEM_MALLOC_HUGE_FIRST_P2P);
aclrtMemcpy(dev1Mem, 10, dev0Mem, 10, ACL_MEMCPY_DEVICE_TO_DEVICE); // 将Device 0上的内存数据复制到Device 1上
aclrtDeviceDisablePeerAccess(0); // 关闭当前Device（Device 1）到指定Device（Device 0）的数据交互
aclrtFree(dev1Mem);
aclrtResetDevice(1); // 释放Device 1的资源
aclrtSetDevice(0); // 切换到Device 0上进行操作
aclrtDeviceDisablePeerAccess(1); // 关闭当前Device（Device 0）到指定Device（Device 1）的数据交互
aclrtFree(dev0Mem);
aclrtResetDeviceForce(0); // Device 0下的操作，调用aclrtResetDeviceForce接口释放Device 0的资源
printf("P2P copy success\n");
} else {
printf("current device doesn't support p2p feature\n");
}
aclFinalize(); // 去初始化